--- gcc/tree-cfg.c.jj	2013-08-27 20:53:36.000000000 +0200
+++ gcc/tree-cfg.c	2013-09-02 13:55:36.591737508 +0200
@@ -610,7 +610,6 @@ make_edges (void)
 	    case GIMPLE_OMP_TASK:
 	    case GIMPLE_OMP_FOR:
 	    case GIMPLE_OMP_SINGLE:
-	    case GIMPLE_OMP_TARGET:
 	    case GIMPLE_OMP_TEAMS:
 	    case GIMPLE_OMP_MASTER:
 	    case GIMPLE_OMP_ORDERED:
@@ -620,6 +619,13 @@ make_edges (void)
 	      fallthru = true;
 	      break;
 
+	    case GIMPLE_OMP_TARGET:
+	      cur_region = new_omp_region (bb, code, cur_region);
+	      fallthru = true;
+	      if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE)
+		cur_region = cur_region->outer;
+	      break;
+
 	    case GIMPLE_OMP_SECTIONS:
 	      cur_region = new_omp_region (bb, code, cur_region);
 	      fallthru = true;
--- gcc/langhooks.c.jj	2013-08-27 20:53:03.000000000 +0200
+++ gcc/langhooks.c	2013-09-04 19:30:55.551423852 +0200
@@ -523,6 +523,15 @@ lhd_omp_firstprivatize_type_sizes (struc
 {
 }
 
+/* Return true if TYPE is an OpenMP mappable type.  By default return true
+   if type is complete.  */
+
+bool
+lhd_omp_mappable_type (tree type)
+{
+  return COMPLETE_TYPE_P (type);
+}
+
 /* Common function for add_builtin_function and
    add_builtin_function_ext_scope.  */
 static tree
--- gcc/omp-low.c.jj	2013-08-27 22:44:31.000000000 +0200
+++ gcc/omp-low.c	2013-09-04 19:58:30.320019227 +0200
@@ -1453,6 +1453,9 @@ scan_sharing_clauses (tree clauses, omp_
 	  break;
 
 	case OMP_CLAUSE_SHARED:
+	  /* Ignore shared directives in teams construct.  */
+	  if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+	    break;
 	  gcc_assert (is_taskreg_ctx (ctx));
 	  decl = OMP_CLAUSE_DECL (c);
 	  gcc_assert (!COMPLETE_TYPE_P (TREE_TYPE (decl))
@@ -1533,6 +1536,9 @@ scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_THREAD_LIMIT:
+	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEPEND:
@@ -1540,6 +1546,66 @@ scan_sharing_clauses (tree clauses, omp_
 	    scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
 	  break;
 
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	case OMP_CLAUSE_MAP:
+	  if (ctx->outer)
+	    scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
+	  decl = OMP_CLAUSE_DECL (c);
+	  /* Global variables with "omp declare target" attribute
+	     don't need to be copied, the receiver side will use them
+	     directly.  */
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && DECL_P (decl)
+	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+	      && lookup_attribute ("omp declare target",
+				   DECL_ATTRIBUTES (decl)))
+	    break;
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && 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
+		 those.  */
+	      if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
+		  && !POINTER_TYPE_P (TREE_TYPE (decl)))
+		break;
+	    }
+	  if (DECL_P (decl))
+	    {
+	      install_var_field (decl, true, 3, ctx);
+	      if (gimple_omp_target_kind (ctx->stmt)
+		  == GF_OMP_TARGET_KIND_REGION)
+		install_var_local (decl, ctx);
+	    }
+	  else
+	    {
+	      tree base = get_base_address (decl);
+	      tree nc = OMP_CLAUSE_CHAIN (c);
+	      if (DECL_P (base)
+		  && nc != NULL_TREE
+		  && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+		  && OMP_CLAUSE_DECL (nc) == base
+		  && OMP_CLAUSE_MAP_KIND (nc) == OMP_CLAUSE_MAP_POINTER
+		  && integer_zerop (OMP_CLAUSE_SIZE (nc)))
+		{
+		  OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) = 1;
+		  OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (nc) = 1;
+		}
+	      else
+		{
+		  gcc_assert (!splay_tree_lookup (ctx->field_map,
+						  (splay_tree_key) decl));
+		  tree field
+		    = build_decl (OMP_CLAUSE_LOCATION (c),
+				  FIELD_DECL, NULL_TREE, ptr_type_node);
+		  insert_field_into_struct (ctx->record_type, field);
+		  splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
+				     (splay_tree_value) field);
+		}
+	    }
+	  break;
+
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_COLLAPSE:
@@ -1590,16 +1656,46 @@ scan_sharing_clauses (tree clauses, omp_
 	  break;
 
 	case OMP_CLAUSE_SHARED:
+	  /* Ignore shared directives in teams construct.  */
+	  if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+	    break;
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (! is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
 	    fixup_remapped_decl (decl, ctx, false);
 	  break;
 
+	case OMP_CLAUSE_MAP:
+	  if (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)))
+	    break;
+	  if (DECL_P (decl))
+	    {
+	      if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+		  && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
+		  && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
+		{
+		  tree new_decl = lookup_decl (decl, ctx);
+		  TREE_TYPE (new_decl)
+		    = remap_type (TREE_TYPE (decl), &ctx->cb);
+		}
+	      else
+		fixup_remapped_decl (decl, ctx, false);
+	    }
+	  break;
+
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_COPYIN:
 	case OMP_CLAUSE_DEFAULT:
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_THREAD_LIMIT:
+	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_NOWAIT:
@@ -1613,6 +1709,8 @@ scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_ALIGNED:
 	case OMP_CLAUSE_DEPEND:
 	case OMP_CLAUSE__LOOPTEMP_:
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
 	  break;
 
 	default:
@@ -1677,6 +1775,26 @@ create_omp_child_function (omp_context *
   DECL_EXTERNAL (decl) = 0;
   DECL_CONTEXT (decl) = NULL_TREE;
   DECL_INITIAL (decl) = make_node (BLOCK);
+  bool target_p = false;
+  if (lookup_attribute ("omp declare target",
+			DECL_ATTRIBUTES (current_function_decl)))
+    target_p = true;
+  else
+    {
+      omp_context *octx;
+      for (octx = ctx; octx; octx = octx->outer)
+	if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
+	    && gimple_omp_target_kind (octx->stmt)
+	       == GF_OMP_TARGET_KIND_REGION)
+	  {
+	    target_p = true;
+	    break;
+	  }
+    }
+  if (target_p)
+    DECL_ATTRIBUTES (decl)
+      = tree_cons (get_identifier ("omp declare target"),
+		   NULL_TREE, DECL_ATTRIBUTES (decl));
 
   t = build_decl (DECL_SOURCE_LOCATION (decl),
 		  RESULT_DECL, NULL_TREE, void_type_node);
@@ -1975,6 +2093,53 @@ scan_omp_single (gimple stmt, omp_contex
     layout_type (ctx->record_type);
 }
 
+/* Scan an OpenMP target{, data, update} directive.  */
+
+static void
+scan_omp_target (gimple stmt, omp_context *outer_ctx)
+{
+  omp_context *ctx;
+  tree name;
+  int kind = gimple_omp_target_kind (stmt);
+
+  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;
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      create_omp_child_function (ctx, false);
+      gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+    }
+
+  scan_sharing_clauses (gimple_omp_target_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
+    {
+      layout_type (ctx->record_type);
+      if (kind == GF_OMP_TARGET_KIND_REGION)
+	fixup_child_record_type (ctx);
+    }
+}
+
+/* Scan an OpenMP teams directive.  */
+
+static void
+scan_omp_teams (gimple stmt, omp_context *outer_ctx)
+{
+  omp_context *ctx = new_omp_context (stmt, outer_ctx);
+  scan_sharing_clauses (gimple_omp_teams_clauses (stmt), ctx);
+  scan_omp (gimple_omp_body_ptr (stmt), ctx);
+}
 
 /* Check OpenMP nesting restrictions.  */
 static bool
@@ -1992,7 +2157,7 @@ check_omp_nesting_restrictions (gimple s
       else if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
 	{
 	  if ((gimple_code (stmt) != GIMPLE_OMP_FOR
-	       || (gimple_omp_for_kind (ctx->stmt)
+	       || (gimple_omp_for_kind (stmt)
 		   != GF_OMP_FOR_KIND_DISTRIBUTE))
 	      && gimple_code (stmt) != GIMPLE_OMP_PARALLEL)
 	    {
@@ -2345,6 +2510,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *g
       scan_omp (gimple_omp_body_ptr (stmt), ctx);
       break;
 
+    case GIMPLE_OMP_TARGET:
+      scan_omp_target (stmt, ctx);
+      break;
+
+    case GIMPLE_OMP_TEAMS:
+      scan_omp_teams (stmt, ctx);
+      break;
+
     case GIMPLE_BIND:
       {
 	tree var;
@@ -2731,6 +2904,9 @@ lower_rec_input_clauses (tree clauses, g
 		continue;
 	      break;
 	    case OMP_CLAUSE_SHARED:
+	      /* Ignore shared directives in teams construct.  */
+	      if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+		continue;
 	      if (maybe_lookup_decl (OMP_CLAUSE_DECL (c), ctx) == NULL)
 		{
 		  gcc_assert (is_global_var (OMP_CLAUSE_DECL (c)));
@@ -2889,6 +3065,9 @@ lower_rec_input_clauses (tree clauses, g
 	  switch (OMP_CLAUSE_CODE (c))
 	    {
 	    case OMP_CLAUSE_SHARED:
+	      /* Ignore shared directives in teams construct.  */
+	      if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+		continue;
 	      /* Shared global vars are just accessed directly.  */
 	      if (is_global_var (new_var))
 		break;
@@ -6728,7 +6907,8 @@ expand_omp_synch (struct omp_region *reg
   gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_SINGLE
 	      || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_MASTER
 	      || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_ORDERED
-	      || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL);
+	      || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL
+	      || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_TEAMS);
   gsi_remove (&si, true);
   single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
 
@@ -7317,6 +7497,318 @@ expand_omp_atomic (struct omp_region *re
 }
 
 
+/* Expand the OpenMP target{, data, update} directive starting at REGION.  */
+
+static void
+expand_omp_target (struct omp_region *region)
+{
+  basic_block entry_bb, exit_bb, new_bb;
+  struct function *child_cfun = NULL;
+  tree child_fn = NULL_TREE, block, t;
+  gimple_stmt_iterator gsi;
+  gimple entry_stmt, stmt;
+  edge e;
+
+  entry_stmt = last_stmt (region->entry);
+  new_bb = region->entry;
+  int kind = gimple_omp_target_kind (entry_stmt);
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      child_fn = gimple_omp_target_child_fn (entry_stmt);
+      child_cfun = DECL_STRUCT_FUNCTION (child_fn);
+    }
+
+  entry_bb = region->entry;
+  exit_bb = region->exit;
+
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      unsigned srcidx, dstidx, num;
+
+      /* If the target 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_omp_target_data_arg (entry_stmt))
+	{
+	  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_target_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 ignoring the subcode because we're
+		     effectively doing a STRIP_NOPS.  */
+
+		  if (TREE_CODE (arg) == ADDR_EXPR
+		      && TREE_OPERAND (arg, 0) == sender)
+		    {
+		      tgtcopy_stmt = stmt;
+		      break;
+		    }
+		}
+	    }
+
+	  gcc_assert (tgtcopy_stmt != NULL);
+	  arg = DECL_ARGUMENTS (child_fn);
+
+	  gcc_assert (gimple_assign_lhs (tgtcopy_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 target 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_OMP_TARGET,
+	 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_OMP_TARGET
+		  && gimple_omp_target_kind (stmt)
+		     == GF_OMP_TARGET_KIND_REGION);
+      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 target 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 OMP 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 the target region, or do data
+     transfers.  */
+  tree t1, t2, t3, t4, device, cond, c, clauses;
+  enum built_in_function start_ix;
+  location_t clause_loc;
+
+  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;
+
+  /* By default, the value of DEVICE is -1 (let runtime library choose)
+     and there is no conditional.  */
+  cond = NULL_TREE;
+  device = build_int_cst (integer_type_node, -1);
+
+  c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+  if (c)
+    cond = OMP_CLAUSE_IF_EXPR (c);
+
+  c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+  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);
+
+  /* If we found the clause 'if (cond)', build
+     (cond ? device : -2).  */
+  if (cond)
+    {
+      cond = gimple_boolify (cond);
+
+      basic_block cond_bb, then_bb, else_bb;
+      edge e;
+      tree tmp_var;
+
+      tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
+      e = split_block (entry_bb, NULL);
+      cond_bb = e->src;
+      entry_bb = e->dest;
+      remove_edge (e);
+
+      then_bb = create_empty_bb (cond_bb);
+      else_bb = create_empty_bb (then_bb);
+      set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+      set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
+
+      stmt = gimple_build_cond_empty (cond);
+      gsi = gsi_start_bb (cond_bb);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      gsi = gsi_start_bb (then_bb);
+      stmt = gimple_build_assign (tmp_var, device);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      gsi = gsi_start_bb (else_bb);
+      stmt = gimple_build_assign (tmp_var,
+				  build_int_cst (integer_type_node, -2));
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+      make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
+      if (current_loops)
+	{
+	  add_bb_to_loop (then_bb, cond_bb->loop_father);
+	  add_bb_to_loop (else_bb, cond_bb->loop_father);
+	}
+      make_edge (then_bb, entry_bb, EDGE_FALLTHRU);
+      make_edge (else_bb, entry_bb, EDGE_FALLTHRU);
+
+      device = tmp_var;
+    }
+
+  gsi = gsi_last_bb (new_bb);
+  t = gimple_omp_target_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;
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      tree fnaddr = build_fold_addr_expr (child_fn);
+      unsigned fnnamelen = IDENTIFIER_LENGTH (DECL_NAME (child_fn));
+      tree fnname = build_string (fnnamelen,
+				  IDENTIFIER_POINTER (DECL_NAME (child_fn)));
+      TREE_TYPE (fnname) = build_array_type_nelts (char_type_node,
+						   fnnamelen);
+      TREE_READONLY (fnname) = 1;
+      TREE_STATIC (fnname) = 1;
+      fnname = build_fold_addr_expr (fnname);
+      g = gimple_build_call (builtin_decl_explicit (start_ix), 7,
+			     device, fnaddr, fnname, t1, t2, t3, t4);
+    }
+  else
+    g = gimple_build_call (builtin_decl_explicit (start_ix), 5,
+			   device, t1, t2, t3, t4);
+  gimple_set_location (g, gimple_location (entry_stmt));
+  gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+  if (kind != GF_OMP_TARGET_KIND_REGION)
+    {
+      g = gsi_stmt (gsi);
+      gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
+      gsi_remove (&gsi, true);
+    }
+  if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+    {
+      gsi = gsi_last_bb (region->exit);
+      g = gsi_stmt (gsi);
+      gcc_assert (g && gimple_code (g) == GIMPLE_OMP_RETURN);
+      gsi_remove (&gsi, true);
+    }
+}
+
+
 /* Expand the parallel region tree rooted at REGION.  Expansion
    proceeds in depth-first order.  Innermost regions are expanded
    first.  This way, parallel regions that require a new function to
@@ -7374,6 +7866,7 @@ expand_omp (struct omp_region *region)
 	case GIMPLE_OMP_MASTER:
 	case GIMPLE_OMP_ORDERED:
 	case GIMPLE_OMP_CRITICAL:
+	case GIMPLE_OMP_TEAMS:
 	  expand_omp_synch (region);
 	  break;
 
@@ -7381,6 +7874,10 @@ expand_omp (struct omp_region *region)
 	  expand_omp_atomic (region);
 	  break;
 
+	case GIMPLE_OMP_TARGET:
+	  expand_omp_target (region);
+	  break;
+
 	default:
 	  gcc_unreachable ();
 	}
@@ -7445,6 +7942,9 @@ build_omp_regions_1 (basic_block bb, str
 	     GIMPLE_OMP_SECTIONS, and we do nothing for it.  */
 	  ;
 	}
+      else if (code == GIMPLE_OMP_TARGET
+	       && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE)
+	new_omp_region (bb, code, parent);
       else
 	{
 	  /* Otherwise, this directive becomes the parent for a new
@@ -7825,7 +8325,7 @@ lower_omp_single (gimple_stmt_iterator *
 {
   tree block;
   gimple t, bind, single_stmt = gsi_stmt (*gsi_p);
-  gimple_seq bind_body, dlist;
+  gimple_seq bind_body, bind_body_tail = NULL, dlist;
   struct gimplify_ctx gctx;
 
   push_gimplify_context (&gctx);
@@ -7855,8 +8355,17 @@ lower_omp_single (gimple_stmt_iterator *
   t = gimple_build_omp_return
         (!!find_omp_clause (gimple_omp_single_clauses (single_stmt),
 			    OMP_CLAUSE_NOWAIT));
-  gimple_seq_add_stmt (&bind_body, t);
-  maybe_add_implicit_barrier_cancel (ctx, &bind_body);
+  gimple_seq_add_stmt (&bind_body_tail, t);
+  maybe_add_implicit_barrier_cancel (ctx, &bind_body_tail);
+  if (ctx->record_type)
+    {
+      gimple_stmt_iterator gsi = gsi_start (bind_body_tail);
+      tree clobber = build_constructor (ctx->record_type, NULL);
+      TREE_THIS_VOLATILE (clobber) = 1;
+      gsi_insert_after (&gsi, gimple_build_assign (ctx->sender_decl,
+						   clobber), GSI_SAME_STMT);
+    }
+  gimple_seq_add_seq (&bind_body, bind_body_tail);
   gimple_bind_set_body (bind, bind_body);
 
   pop_gimplify_context (bind);
@@ -8605,6 +9114,14 @@ lower_omp_taskreg (gimple_stmt_iterator
   lower_send_clauses (clauses, &ilist, &olist, ctx);
   lower_send_shared_vars (&ilist, &olist, ctx);
 
+  if (ctx->record_type)
+    {
+      tree clobber = build_constructor (TREE_TYPE (ctx->sender_decl), 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.  */
 
@@ -8637,6 +9154,326 @@ lower_omp_taskreg (gimple_stmt_iterator
   pop_gimplify_context (NULL);
 }
 
+/* Lower the OpenMP target directive in the current statement
+   in GSI_P.  CTX holds context information for the directive.  */
+
+static void
+lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+  tree clauses;
+  tree child_fn, t, c;
+  gimple stmt = gsi_stmt (*gsi_p);
+  gimple tgt_bind = NULL, bind;
+  gimple_seq tgt_body = NULL, olist, ilist, new_body;
+  struct gimplify_ctx gctx;
+  location_t loc = gimple_location (stmt);
+  int kind = gimple_omp_target_kind (stmt);
+  unsigned int map_cnt = 0;
+
+  clauses = gimple_omp_target_clauses (stmt);
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      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)
+    tgt_body = gimple_omp_body (stmt);
+  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 (!lookup_sfield (var, ctx))
+	  continue;
+
+	if (kind == GF_OMP_TARGET_KIND_REGION)
+	  {
+	    x = build_receiver_ref (var, true, ctx);
+	    tree new_var = lookup_decl (var, ctx);
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	map_cnt++;
+      }
+
+  if (kind != GF_OMP_TARGET_KIND_UPDATE)
+    lower_omp (&tgt_body, ctx);
+
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      /* Declare all the variables created by mapping and the variables
+	 declared in the scope of the target body.  */
+      record_vars_into (ctx->block_vars, child_fn);
+      record_vars_into (gimple_bind_vars (tgt_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_omp_target_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 (!lookup_sfield (ovar, ctx))
+	      continue;
+
+	    if (nc)
+	      {
+		tree var = lookup_decl_in_outer_ctx (ovar, ctx);
+		tree x = build_sender_ref (ovar, ctx);
+		if (is_gimple_reg (var))
+		  {
+		    gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
+		    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 (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 ();
+	      }
+	    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 && kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      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));
+    }
+
+  if (kind == GF_OMP_TARGET_KIND_REGION)
+    {
+      gimple_seq_add_seq (&new_body, tgt_body);
+      new_body = maybe_catch_exception (new_body);
+    }
+  else if (kind == GF_OMP_TARGET_KIND_DATA)
+    new_body = tgt_body;
+  if (kind != GF_OMP_TARGET_KIND_UPDATE)
+    {
+      gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
+      gimple_omp_set_body (stmt, new_body);
+    }
+
+  bind = gimple_build_bind (NULL, NULL,
+			    tgt_bind ? gimple_bind_block (tgt_bind)
+				     : NULL_TREE);
+  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);
+}
+
+/* Expand code for an OpenMP teams directive.  */
+
+static void
+lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+  gimple teams_stmt = gsi_stmt (*gsi_p);
+  struct gimplify_ctx gctx;
+  push_gimplify_context (&gctx);
+
+  tree block = make_node (BLOCK);
+  gimple bind = gimple_build_bind (NULL, NULL, block);
+  gsi_replace (gsi_p, bind, true);
+  gimple_seq bind_body = NULL;
+  gimple_seq dlist = NULL;
+  gimple_seq olist = NULL;
+
+  tree num_teams = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
+				    OMP_CLAUSE_NUM_TEAMS);
+  if (num_teams == NULL_TREE)
+    num_teams = build_int_cst (unsigned_type_node, 0);
+  else
+    {
+      num_teams = OMP_CLAUSE_NUM_TEAMS_EXPR (num_teams);
+      num_teams = fold_convert (unsigned_type_node, num_teams);
+      gimplify_expr (&num_teams, &bind_body, NULL, is_gimple_val, fb_rvalue);
+    }
+  tree thread_limit = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
+				       OMP_CLAUSE_THREAD_LIMIT);
+  if (thread_limit == NULL_TREE)
+    thread_limit = build_int_cst (unsigned_type_node, 0);
+  else
+    {
+      thread_limit = OMP_CLAUSE_THREAD_LIMIT_EXPR (thread_limit);
+      thread_limit = fold_convert (unsigned_type_node, thread_limit);
+      gimplify_expr (&thread_limit, &bind_body, NULL, is_gimple_val,
+		     fb_rvalue);
+    }
+
+  lower_rec_input_clauses (gimple_omp_teams_clauses (teams_stmt),
+			   &bind_body, &dlist, ctx, NULL);
+  lower_omp (gimple_omp_body_ptr (teams_stmt), ctx);
+  lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist, ctx);
+  gimple_seq_add_stmt (&bind_body, teams_stmt);
+
+  location_t loc = gimple_location (teams_stmt);
+  tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS);
+  gimple call = gimple_build_call (decl, 2, num_teams, thread_limit);
+  gimple_set_location (call, loc);
+  gimple_seq_add_stmt (&bind_body, call);
+
+  gimple_seq_add_seq (&bind_body, gimple_omp_body (teams_stmt));
+  gimple_omp_set_body (teams_stmt, NULL);
+  gimple_seq_add_seq (&bind_body, olist);
+  gimple_seq_add_seq (&bind_body, dlist);
+  gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true));
+  gimple_bind_set_body (bind, bind_body);
+
+  pop_gimplify_context (bind);
+
+  gimple_bind_append_vars (bind, ctx->block_vars);
+  BLOCK_VARS (block) = ctx->block_vars;
+  if (BLOCK_VARS (block))
+    TREE_USED (block) = 1;
+}
+
+
 /* Callback for lower_omp_1.  Return non-NULL if *tp needs to be
    regimplified.  If DATA is non-NULL, lower_omp_1 is outside
    of OpenMP context, but with task_shared_vars set.  */
@@ -8760,6 +9597,16 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p
 			lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
 	gimple_regimplify_operands (stmt, gsi_p);
       break;
+    case GIMPLE_OMP_TARGET:
+      ctx = maybe_lookup_ctx (stmt);
+      gcc_assert (ctx);
+      lower_omp_target (gsi_p, ctx);
+      break;
+    case GIMPLE_OMP_TEAMS:
+      ctx = maybe_lookup_ctx (stmt);
+      gcc_assert (ctx);
+      lower_omp_teams (gsi_p, ctx);
+      break;
     case GIMPLE_CALL:
       tree fndecl;
       fndecl = gimple_call_fndecl (stmt);
--- gcc/c/c-typeck.c.jj	2013-08-27 20:49:05.000000000 +0200
+++ gcc/c/c-typeck.c	2013-09-04 17:21:17.295112407 +0200
@@ -10773,8 +10773,7 @@ c_finish_omp_cancellation_point (locatio
 
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
-			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     bool &pointer_based_p)
+			     bool &maybe_zero_len, unsigned int &first_non_one)
 {
   tree ret, low_bound, length, type;
   if (TREE_CODE (t) != TREE_LIST)
@@ -10801,15 +10800,11 @@ handle_omp_array_sections_1 (tree c, tre
 		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	  return error_mark_node;
 	}
-      if (POINTER_TYPE_P (TREE_TYPE (t))
-	  && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
-	pointer_based_p = true;
       return t;
     }
 
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
-				     maybe_zero_len, first_non_one,
-				     pointer_based_p);
+				     maybe_zero_len, first_non_one);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -10989,16 +10984,12 @@ handle_omp_array_sections_1 (tree c, tre
     }
   if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
     types.safe_push (TREE_TYPE (ret));
-  /* For pointer based array sections we will need to evaluate lb more
-     than once.  */
-  if (pointer_based_p)
+  /* We will need to evaluate lb more than once.  */
+  tree lb = c_save_expr (low_bound);
+  if (lb != low_bound)
     {
-      tree lb = c_save_expr (low_bound);
-      if (lb != low_bound)
-	{
-	  TREE_PURPOSE (t) = lb;
-	  low_bound = lb;
-	}
+      TREE_PURPOSE (t) = lb;
+      low_bound = lb;
     }
   ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
   return ret;
@@ -11010,12 +11001,10 @@ static bool
 handle_omp_array_sections (tree c)
 {
   bool maybe_zero_len = false;
-  bool pointer_based_p = false;
   unsigned int first_non_one = 0;
   vec<tree> types = vNULL;
   tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
-					    maybe_zero_len, first_non_one,
-					    pointer_based_p);
+					    maybe_zero_len, first_non_one);
   if (first == error_mark_node)
     {
       types.release ();
@@ -11047,6 +11036,7 @@ handle_omp_array_sections (tree c)
 	}
       if (tem)
 	first = build2 (COMPOUND_EXPR, TREE_TYPE (first), tem, first);
+      first = c_fully_fold (first, false, NULL);
       OMP_CLAUSE_DECL (c) = first;
     }
   else
@@ -11155,27 +11145,31 @@ handle_omp_array_sections (tree c)
       types.release ();
       if (side_effects)
 	size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
+      first = c_fully_fold (first, false, NULL);
       OMP_CLAUSE_DECL (c) = first;
+      if (size)
+	size = c_fully_fold (size, false, NULL);
       OMP_CLAUSE_SIZE (c) = size;
-      if (pointer_based_p)
-	{
-	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
-	  if (!c_mark_addressable (t))
-	    return false;
-	  OMP_CLAUSE_DECL (c2) = t;
-	  t = build_fold_addr_expr (first);
-	  t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
-				build_pointer_type (char_type_node), t);
-	  t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
-			       ptrdiff_type_node, t,
-			       fold_convert_loc (OMP_CLAUSE_LOCATION (c),
-						 TREE_TYPE (t),
-						 OMP_CLAUSE_DECL (c2)));
-	  OMP_CLAUSE_SIZE (c2) = t;
-	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
-	  OMP_CLAUSE_CHAIN (c) = c2;
-	}
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+	return false;
+      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+      OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
+      if (!c_mark_addressable (t))
+	return false;
+      OMP_CLAUSE_DECL (c2) = t;
+      t = build_fold_addr_expr (first);
+      t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
+      tree ptr = OMP_CLAUSE_DECL (c2);
+      if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
+	ptr = build_fold_addr_expr (ptr);
+      t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+			   ptrdiff_type_node, t,
+			   fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+					     ptrdiff_type_node, ptr));
+      t = c_fully_fold (t, false, NULL);
+      OMP_CLAUSE_SIZE (c2) = t;
+      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+      OMP_CLAUSE_CHAIN (c) = c2;
     }
   return false;
 }
@@ -11462,18 +11456,21 @@ c_finish_omp_clauses (tree clauses)
 	    }
 	  else if (!c_mark_addressable (t))
 	    remove = true;
-	  else if (!COMPLETE_TYPE_P (TREE_TYPE (t)))
+	  else if (!COMPLETE_TYPE_P (TREE_TYPE (t))
+		   && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qD does not have a mappable type in %qs clause", t,
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
-	    break;
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
 	    {
-	      error ("%qD appears more than once in motion clauses", t);
+	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+		error ("%qD appears more than once in motion clauses", t);
+	      else
+		error ("%qD appears more than once in map clauses", t);
 	      remove = true;
 	    }
 	  else
@@ -11509,6 +11506,7 @@ c_finish_omp_clauses (tree clauses)
 
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
+	case OMP_CLAUSE_NUM_TEAMS:
 	case OMP_CLAUSE_THREAD_LIMIT:
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_ORDERED:
--- gcc/c/c-parser.c.jj	2013-08-27 20:49:04.000000000 +0200
+++ gcc/c/c-parser.c	2013-09-02 11:41:51.023319550 +0200
@@ -12068,7 +12068,10 @@ c_parser_omp_target_data (location_t loc
   OMP_TARGET_DATA_CLAUSES (stmt)
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				"#pragma omp target data");
-  OMP_TARGET_DATA_BODY (stmt) = c_parser_omp_structured_block (parser);
+  keep_next_level ();
+  tree block = c_begin_compound_stmt (true);
+  add_stmt (c_parser_omp_structured_block (parser));
+  OMP_TARGET_DATA_BODY (stmt) = c_end_compound_stmt (loc, block, true);
 
   SET_EXPR_LOCATION (stmt, loc);
   return add_stmt (stmt);
@@ -12161,6 +12164,7 @@ c_parser_omp_target (c_parser *parser, e
 
 	  c_parser_consume_token (parser);
 	  strcpy (p_name, "#pragma omp target");
+	  keep_next_level ();
 	  tree block = c_begin_compound_stmt (true);
 	  tree ret = c_parser_omp_teams (loc, parser, p_name,
 					 OMP_TARGET_CLAUSE_MASK, cclauses);
@@ -12182,7 +12186,10 @@ c_parser_omp_target (c_parser *parser, e
   OMP_TARGET_CLAUSES (stmt)
     = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				"#pragma omp target");
-  OMP_TARGET_BODY (stmt) = c_parser_omp_structured_block (parser);
+  keep_next_level ();
+  tree block = c_begin_compound_stmt (true);
+  add_stmt (c_parser_omp_structured_block (parser));
+  OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
 
   SET_EXPR_LOCATION (stmt, loc);
   add_stmt (stmt);
--- gcc/tree.h.jj	2013-08-27 22:05:43.000000000 +0200
+++ gcc/tree.h	2013-09-03 14:35:36.350942556 +0200
@@ -629,6 +629,9 @@ struct GTY(()) tree_base {
        OMP_CLAUSE_LINEAR_NO_COPYIN in
 	   OMP_CLAUSE_LINEAR
 
+       OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION in
+	   OMP_CLAUSE_MAP
+
        TRANSACTION_EXPR_RELAXED in
 	   TRANSACTION_EXPR
 
@@ -2050,15 +2053,21 @@ enum omp_clause_map_kind
   OMP_CLAUSE_MAP_TO,
   OMP_CLAUSE_MAP_FROM,
   OMP_CLAUSE_MAP_TOFROM,
-  /* This following is an internal only map kind, used for pointer based array
-     sections.  OMP_CLAUSE_SIZE for these is not the pointer size, which is
-     implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias.  */
+  /* The following kind is an internal only map kind, used for pointer based
+     array sections.  OMP_CLAUSE_SIZE for these is not the pointer size,
+     which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias.  */
   OMP_CLAUSE_MAP_POINTER
 };
 
 #define OMP_CLAUSE_MAP_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
 
+/* Nonzero if this map clause is for array (rather than pointer) based array
+   section with zero bias.  Both the non-decl OMP_CLAUSE_MAP and
+   correspoidng OMP_CLAUSE_MAP_POINTER clause are marked with this flag.  */
+#define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag)
+
 enum omp_clause_proc_bind_kind
 {
   /* Numbers should match omp_proc_bind_t enum in omp.h.  */
--- gcc/gimple.def.jj	2013-08-19 12:07:51.000000000 +0200
+++ gcc/gimple.def	2013-09-02 11:55:57.892684945 +0200
@@ -349,12 +349,18 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "g
    CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
 DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE)
 
-/* GIMPLE_OMP_TARGET <BODY, CLAUSES> represents
+/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
    #pragma omp target {,data,update}
    BODY is the sequence of statements inside the target construct
    (NULL for target update).
-   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
-DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE)
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+   CHILD_FN is set when outlining the body of the target 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_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL)
 
 /* GIMPLE_OMP_TEAMS <BODY, CLAUSES> represents #pragma omp teams
    BODY is the sequence of statements inside the single section.
--- gcc/fortran/types.def.jj	2013-08-19 12:07:51.000000000 +0200
+++ gcc/fortran/types.def	2013-09-04 20:19:40.880566458 +0200
@@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
 DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
 		     BT_CONST_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
 
@@ -167,6 +168,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_LONG_LON
 		     BT_PTR_LONG, BT_PTR_LONG)
 DEF_FUNCTION_TYPE_5 (BT_FN_VOID_SIZE_VPTR_PTR_PTR_INT, BT_VOID, BT_SIZE,
 		     BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
                      BT_BOOL, BT_LONG, BT_LONG, BT_LONG, BT_LONG,
@@ -203,6 +206,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
 		     BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
 		     BT_ULONGLONG, BT_ULONGLONG,
 		     BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+		     BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
--- gcc/cp/parser.c.jj	2013-08-27 20:50:58.000000000 +0200
+++ gcc/cp/parser.c	2013-09-04 17:37:20.511076432 +0200
@@ -29451,6 +29451,7 @@ cp_parser_omp_target_data (cp_parser *pa
   OMP_TARGET_DATA_CLAUSES (stmt)
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				 "#pragma omp target data", pragma_tok);
+  keep_next_level (true);
   OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser);
 
   SET_EXPR_LOCATION (stmt, pragma_tok->location);
@@ -29543,6 +29544,7 @@ cp_parser_omp_target (cp_parser *parser,
 
 	  cp_lexer_consume_token (parser->lexer);
 	  strcpy (p_name, "#pragma omp target");
+	  keep_next_level (true);
 	  tree sb = begin_omp_structured_block ();
 	  unsigned save = cp_parser_begin_omp_structured_block (parser);
 	  tree ret = cp_parser_omp_teams (parser, pragma_tok, p_name,
@@ -29566,6 +29568,7 @@ cp_parser_omp_target (cp_parser *parser,
   OMP_TARGET_CLAUSES (stmt)
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				 "#pragma omp target", pragma_tok);
+  keep_next_level (true);
   OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
 
   SET_EXPR_LOCATION (stmt, pragma_tok->location);
--- gcc/cp/cp-objcp-common.h.jj	2013-03-20 10:07:19.000000000 +0100
+++ gcc/cp/cp-objcp-common.h	2013-09-04 19:33:40.520583302 +0200
@@ -145,6 +145,8 @@ extern void cp_common_init_ts (void);
 #define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
 #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
 #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
+#undef LANG_HOOKS_OMP_MAPPABLE_TYPE
+#define LANG_HOOKS_OMP_MAPPABLE_TYPE cp_omp_mappable_type
 
 #undef LANG_HOOKS_EH_USE_CXA_END_CLEANUP
 #define LANG_HOOKS_EH_USE_CXA_END_CLEANUP true
--- gcc/cp/semantics.c.jj	2013-08-27 20:50:56.000000000 +0200
+++ gcc/cp/semantics.c	2013-09-04 17:21:48.507946784 +0200
@@ -4101,8 +4101,7 @@ cxx_omp_create_clause_info (tree c, tree
 
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
-			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     bool &pointer_based_p)
+			     bool &maybe_zero_len, unsigned int &first_non_one)
 {
   tree ret, low_bound, length, type;
   if (TREE_CODE (t) != TREE_LIST)
@@ -4134,16 +4133,11 @@ handle_omp_array_sections_1 (tree c, tre
 	  return error_mark_node;
 	}
       t = convert_from_reference (t);
-      if (!processing_template_decl
-	  && POINTER_TYPE_P (TREE_TYPE (t))
-	  && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
-	pointer_based_p = true;
       return t;
     }
 
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
-				     maybe_zero_len, first_non_one,
-				     pointer_based_p);
+				     maybe_zero_len, first_non_one);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -4326,16 +4320,12 @@ handle_omp_array_sections_1 (tree c, tre
     }
   if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
     types.safe_push (TREE_TYPE (ret));
-  /* For pointer based array sections we will need to evaluate lb more
-     than once.  */
-  if (pointer_based_p)
+  /* We will need to evaluate lb more than once.  */
+  tree lb = cp_save_expr (low_bound);
+  if (lb != low_bound)
     {
-      tree lb = cp_save_expr (low_bound);
-      if (lb != low_bound)
-	{
-	  TREE_PURPOSE (t) = lb;
-	  low_bound = lb;
-	}
+      TREE_PURPOSE (t) = lb;
+      low_bound = lb;
     }
   ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, false);
   return ret;
@@ -4347,12 +4337,10 @@ static bool
 handle_omp_array_sections (tree c)
 {
   bool maybe_zero_len = false;
-  bool pointer_based_p = false;
   unsigned int first_non_one = 0;
   vec<tree> types = vNULL;
   tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
-					    maybe_zero_len, first_non_one,
-					    pointer_based_p);
+					    maybe_zero_len, first_non_one);
   if (first == error_mark_node)
     {
       types.release ();
@@ -4506,26 +4494,27 @@ handle_omp_array_sections (tree c)
 	    size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
 	  OMP_CLAUSE_DECL (c) = first;
 	  OMP_CLAUSE_SIZE (c) = size;
-	  if (pointer_based_p)
-	    {
-	      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-					  OMP_CLAUSE_MAP);
-	      OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
-	      if (!cxx_mark_addressable (t))
-		return false;
-	      OMP_CLAUSE_DECL (c2) = t;
-	      t = build_fold_addr_expr (first);
-	      t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
-				    build_pointer_type (char_type_node), t);
-	      t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
-				   ptrdiff_type_node, t,
-				   fold_convert_loc (OMP_CLAUSE_LOCATION (c),
-						     TREE_TYPE (t),
-						     OMP_CLAUSE_DECL (c2)));
-	      OMP_CLAUSE_SIZE (c2) = t;
-	      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
-	      OMP_CLAUSE_CHAIN (c) = c2;
-	    }
+	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+	    return false;
+	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
+	  if (!cxx_mark_addressable (t))
+	    return false;
+	  OMP_CLAUSE_DECL (c2) = t;
+	  t = build_fold_addr_expr (first);
+	  t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+				ptrdiff_type_node, t);
+	  tree ptr = OMP_CLAUSE_DECL (c2);
+	  if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
+	    ptr = build_fold_addr_expr (ptr);
+	  t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+			       ptrdiff_type_node, t,
+			       fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+						 ptrdiff_type_node, ptr));
+	  OMP_CLAUSE_SIZE (c2) = t;
+	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+	  OMP_CLAUSE_CHAIN (c) = c2;
 	}
     }
   return false;
@@ -4972,21 +4961,24 @@ finish_omp_clauses (tree clauses)
 		   && TREE_CODE (TREE_TYPE (t)) != REFERENCE_TYPE
 		   && !cxx_mark_addressable (t))
 	    remove = true;
-	  else if (!cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
-					   == REFERENCE_TYPE)
-					  ? TREE_TYPE (TREE_TYPE (t))
-					  : TREE_TYPE (t)))
+	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		     && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+		   && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
+					      == REFERENCE_TYPE)
+					     ? TREE_TYPE (TREE_TYPE (t))
+					     : TREE_TYPE (t)))
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%qD does not have a mappable type in %qs clause", t,
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
-	    break;
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
 	    {
-	      error ("%qD appears more than once in motion clauses", t);
+	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+		error ("%qD appears more than once in motion clauses", t);
+	      else
+		error ("%qD appears more than once in map clauses", t);
 	      remove = true;
 	    }
 	  else
--- gcc/langhooks.h.jj	2013-03-20 10:08:13.000000000 +0100
+++ gcc/langhooks.h	2013-09-04 19:29:20.429909391 +0200
@@ -111,6 +111,9 @@ struct lang_hooks_for_types
      firstprivate variables.  */
   void (*omp_firstprivatize_type_sizes) (struct gimplify_omp_ctx *, tree);
 
+  /* Return true if TYPE is a mappable type.  */
+  bool (*omp_mappable_type) (tree type);
+
   /* Return TRUE if TYPE1 and TYPE2 are identical for type hashing purposes.
      Called only after doing all language independent checks.
      At present, this function is only called when both TYPE1 and TYPE2 are
--- gcc/langhooks-def.h.jj	2013-03-20 10:08:41.000000000 +0100
+++ gcc/langhooks-def.h	2013-09-04 19:29:45.450781052 +0200
@@ -77,6 +77,7 @@ extern tree lhd_omp_assignment (tree, tr
 struct gimplify_omp_ctx;
 extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
+extern bool lhd_omp_mappable_type (tree);
 
 #define LANG_HOOKS_NAME			"GNU unknown"
 #define LANG_HOOKS_IDENTIFIER_SIZE	sizeof (struct lang_identifier)
@@ -166,6 +167,7 @@ extern tree lhd_make_node (enum tree_cod
 #define LANG_HOOKS_TYPE_MAX_SIZE	lhd_return_null_const_tree
 #define LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES \
   lhd_omp_firstprivatize_type_sizes
+#define LANG_HOOKS_OMP_MAPPABLE_TYPE	lhd_omp_mappable_type
 #define LANG_HOOKS_TYPE_HASH_EQ		NULL
 #define LANG_HOOKS_GET_ARRAY_DESCR_INFO	NULL
 #define LANG_HOOKS_GET_SUBRANGE_BOUNDS	NULL
@@ -184,6 +186,7 @@ extern tree lhd_make_node (enum tree_cod
   LANG_HOOKS_INCOMPLETE_TYPE_ERROR, \
   LANG_HOOKS_TYPE_MAX_SIZE, \
   LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES, \
+  LANG_HOOKS_OMP_MAPPABLE_TYPE, \
   LANG_HOOKS_TYPE_HASH_EQ, \
   LANG_HOOKS_GET_ARRAY_DESCR_INFO, \
   LANG_HOOKS_GET_SUBRANGE_BOUNDS, \
--- gcc/gimple.h.jj	2013-08-27 21:55:11.000000000 +0200
+++ gcc/gimple.h	2013-09-02 11:57:32.143164024 +0200
@@ -4702,7 +4702,7 @@ static inline tree
 gimple_omp_target_clauses (const_gimple gs)
 {
   GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
-  return gs->gimple_omp_single.clauses;
+  return gs->gimple_omp_parallel.clauses;
 }
 
 
@@ -4712,7 +4712,7 @@ static inline tree *
 gimple_omp_target_clauses_ptr (gimple gs)
 {
   GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
-  return &gs->gimple_omp_single.clauses;
+  return &gs->gimple_omp_parallel.clauses;
 }
 
 
@@ -4722,7 +4722,7 @@ static inline void
 gimple_omp_target_set_clauses (gimple gs, tree clauses)
 {
   GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
-  gs->gimple_omp_single.clauses = clauses;
+  gs->gimple_omp_parallel.clauses = clauses;
 }
 
 
@@ -4747,6 +4747,67 @@ gimple_omp_target_set_kind (gimple g, in
 }
 
 
+/* Return the child function used to hold the body of OMP_TARGET GS.  */
+
+static inline tree
+gimple_omp_target_child_fn (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return gs->gimple_omp_parallel.child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of
+   OMP_TARGET GS.  */
+
+static inline tree *
+gimple_omp_target_child_fn_ptr (gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return &gs->gimple_omp_parallel.child_fn;
+}
+
+
+/* Set CHILD_FN to be the child function for OMP_TARGET GS.  */
+
+static inline void
+gimple_omp_target_set_child_fn (gimple gs, tree child_fn)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  gs->gimple_omp_parallel.child_fn = child_fn;
+}
+
+
+/* Return the artificial argument used to send variables and values
+   from the parent to the children threads in OMP_TARGET GS.  */
+
+static inline tree
+gimple_omp_target_data_arg (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return gs->gimple_omp_parallel.data_arg;
+}
+
+
+/* Return a pointer to the data argument for OMP_TARGET GS.  */
+
+static inline tree *
+gimple_omp_target_data_arg_ptr (gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return &gs->gimple_omp_parallel.data_arg;
+}
+
+
+/* Set DATA_ARG to be the data argument for OMP_TARGET GS.  */
+
+static inline void
+gimple_omp_target_set_data_arg (gimple gs, tree data_arg)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  gs->gimple_omp_parallel.data_arg = data_arg;
+}
+
+
 /* Return the clauses associated with OMP_TEAMS GS.  */
 
 static inline tree
--- gcc/omp-builtins.def.jj	2013-08-19 12:07:51.000000000 +0200
+++ gcc/omp-builtins.def	2013-09-04 15:14:21.049713033 +0200
@@ -225,3 +225,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C
 		  BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
 		  BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target",
+		  BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+		  ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data",
+		  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
+		  BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
+		  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
+		  BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
--- gcc/gimple-pretty-print.c.jj	2013-08-27 21:53:53.000000000 +0200
+++ gcc/gimple-pretty-print.c	2013-08-30 17:53:14.201347484 +0200
@@ -1277,6 +1277,13 @@ dump_gimple_omp_target (pretty_printer *
       pp_string (buffer, "#pragma omp target");
       pp_string (buffer, kind);
       dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
+      if (gimple_omp_target_child_fn (gs))
+	{
+	  pp_string (buffer, " [child fn: ");
+	  dump_generic_node (buffer, gimple_omp_target_child_fn (gs),
+			     spc, flags, false);
+	  pp_right_bracket (buffer);
+	}
       if (!gimple_seq_empty_p (gimple_omp_body (gs)))
 	{
 	  newline_and_indent (buffer, spc + 2);
--- gcc/gimplify.c.jj	2013-08-27 22:03:35.000000000 +0200
+++ gcc/gimplify.c	2013-09-04 19:36:07.442834807 +0200
@@ -5976,7 +5976,16 @@ omp_notice_variable (struct gimplify_omp
   if (ctx->region_type == ORT_TARGET)
     {
       if (n == NULL)
-	omp_add_variable (ctx, decl, GOVD_MAP | flags);
+	{
+	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+	    {
+	      error ("%qD referenced in target region does not have "
+		     "a mappable type", decl);
+	      omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
+	    }
+	  else
+	    omp_add_variable (ctx, decl, GOVD_MAP | flags);
+	}
       else
 	n->value |= flags;
       ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
@@ -6283,7 +6292,6 @@ gimplify_scan_omp_clauses (tree *list_p,
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
-	  notice_outer = false;
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
@@ -6500,31 +6508,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
   if (private_debug)
     code = OMP_CLAUSE_PRIVATE;
   else if (flags & GOVD_MAP)
-    {
-      /* If decl is already in the enclosing device data environment,
-	 the spec says that it should just be used and no init/assignment
-	 should be done.  If there was any privatization in between though,
-	 it means that original decl might be in the enclosing device data
-	 environment, but the privatized might not.  */
-      struct gimplify_omp_ctx *ctx;
-      for (ctx = gimplify_omp_ctxp->outer_context;
-	   ctx; ctx = ctx->outer_context)
-	{
-	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
-	  if (n == NULL)
-	    continue;
-	  if (ctx->region_type == ORT_TARGET_DATA)
-	    {
-	      if ((n->value & GOVD_MAP) != 0)
-		return 0;
-	    }
-	  else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
-				| GOVD_PRIVATE | GOVD_REDUCTION
-				| GOVD_LINEAR)) != 0)
-	    break;
-	}
-      code = OMP_CLAUSE_MAP;
-    }
+    code = OMP_CLAUSE_MAP;
   else if (flags & GOVD_SHARED)
     {
       if (is_global_var (decl))
@@ -6689,31 +6673,6 @@ gimplify_adjust_omp_clauses (tree *list_
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
 	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
 	    remove = true;
-	  else
-	    {
-	      /* If decl is already in the enclosing device data environment,
-		 the spec says that it should just be used and no init/assignment
-		 should be done.  If there was any privatization in between though,
-		 it means that original decl might be in the enclosing device data
-		 environment, but the privatized might not.  */
-	      struct gimplify_omp_ctx *octx;
-	      for (octx = ctx->outer_context; octx; octx = octx->outer_context)
-		{
-		  n = splay_tree_lookup (octx->variables,
-					 (splay_tree_key) decl);
-		  if (n == NULL)
-		    continue;
-		  if (octx->region_type == ORT_TARGET_DATA)
-		    {
-		      if ((n->value & GOVD_MAP) != 0)
-			remove = true;
-		    }
-		  else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
-					| GOVD_PRIVATE | GOVD_REDUCTION
-					| GOVD_LINEAR)) != 0)
-		    break;
-		}
-	    }
 	  break;
 
 	case OMP_CLAUSE_REDUCTION:
@@ -6722,6 +6681,7 @@ gimplify_adjust_omp_clauses (tree *list_
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_THREAD_LIMIT:
 	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_SCHEDULE:
@@ -7198,7 +7158,28 @@ gimplify_omp_workshare (tree *expr_p, gi
       gcc_unreachable ();
     }
   gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
-  gimplify_and_add (OMP_BODY (expr), &body);
+  if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+    {
+      struct gimplify_ctx gctx;
+      push_gimplify_context (&gctx);
+      gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
+      if (gimple_code (g) == GIMPLE_BIND)
+	pop_gimplify_context (g);
+      else
+	pop_gimplify_context (NULL);
+      if (ort == ORT_TARGET_DATA)
+	{
+	  gimple_seq cleanup = NULL;
+	  tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+	  g = gimple_build_call (fn, 0);
+	  gimple_seq_add_stmt (&cleanup, g);
+	  g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
+	  body = NULL;
+	  gimple_seq_add_stmt (&body, g);
+	}
+    }
+  else
+    gimplify_and_add (OMP_BODY (expr), &body);
   gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
 
   switch (TREE_CODE (expr))
--- gcc/testsuite/c-c++-common/gomp/map-1.c.jj	2013-07-06 19:18:46.161478820 +0200
+++ gcc/testsuite/c-c++-common/gomp/map-1.c	2013-09-04 17:00:51.174322082 +0200
@@ -33,7 +33,7 @@ foo (int g[3][10], int h[4][8], int i[2]
   #pragma omp target map(to: o[2:5]) /* { dg-error "does not have pointer or array type" } */
     ;
   #pragma omp target map(to: a[:][:]) /* { dg-error "array type length expression is not optional" } */
-    bar (&a[0][0]);
+    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" } */
   #pragma omp target map(tofrom: b[-1:]) /* { dg-error "negative low bound in array section" } */
     bar (b);
   #pragma omp target map(tofrom: c[:-3][:]) /* { dg-error "negative length in array section" } */
--- gcc/builtin-types.def.jj	2013-08-19 12:07:51.000000000 +0200
+++ gcc/builtin-types.def	2013-09-03 08:58:48.902658867 +0200
@@ -345,6 +345,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
 DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
 		     BT_CONST_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
 
 DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
 
@@ -472,6 +473,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_VPTR_PTR
 DEF_FUNCTION_TYPE_5 (BT_FN_VOID_OMPFN_PTR_UINT_UINT_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT,
 		     BT_UINT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_6 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VALIST_ARG,
 		     BT_INT, BT_STRING, BT_SIZE, BT_INT, BT_SIZE,
@@ -511,6 +514,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
 		     BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
 		     BT_ULONGLONG, BT_ULONGLONG,
 		     BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+		     BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
--- libgomp/libgomp.map.jj	2013-06-12 14:24:42.000000000 +0200
+++ libgomp/libgomp.map	2013-09-04 15:17:42.694659418 +0200
@@ -218,4 +218,9 @@ GOMP_4.0 {
 	GOMP_parallel;
 	GOMP_taskgroup_start;
 	GOMP_taskgroup_end;
+	GOMP_target;
+	GOMP_target_data;
+	GOMP_target_end_data;
+	GOMP_target_update;
+	GOMP_teams;
 } GOMP_3.0;
--- libgomp/Makefile.am.jj	2013-03-20 10:02:06.000000000 +0100
+++ libgomp/Makefile.am	2013-09-04 14:59:19.475033836 +0200
@@ -60,7 +60,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_L
 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
+	time.c fortran.c affinity.c target.c
 
 nodist_noinst_HEADERS = libgomp_f.h
 nodist_libsubinclude_HEADERS = omp.h
--- libgomp/Makefile.in.jj	2013-03-20 10:02:05.000000000 +0100
+++ libgomp/Makefile.in	2013-09-04 14:59:52.243864494 +0200
@@ -96,7 +96,7 @@ am_libgomp_la_OBJECTS = alloc.lo barrier
 	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
+	fortran.lo affinity.lo target.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_L
 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
+	time.c fortran.c affinity.c target.c
 
 nodist_noinst_HEADERS = libgomp_f.h
 nodist_libsubinclude_HEADERS = omp.h
@@ -474,6 +474,7 @@ distclean-compile:
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sections.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/time.Plo@am__quote@
--- libgomp/libgomp_g.h.jj	2013-04-05 18:15:57.000000000 +0200
+++ libgomp/libgomp_g.h	2013-09-04 15:25:55.982080005 +0200
@@ -199,4 +199,14 @@ extern bool GOMP_single_start (void);
 extern void *GOMP_single_copy_start (void);
 extern void GOMP_single_copy_end (void *);
 
+/* target.c */
+
+extern void GOMP_target (int, void (*) (void *), const char *,
+			 size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_data (int, size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_end_data (void);
+extern void GOMP_target_update (int, size_t, void **, size_t *,
+				unsigned char *);
+extern void GOMP_teams (unsigned int, unsigned int);
+
 #endif /* LIBGOMP_G_H */
--- libgomp/testsuite/libgomp.c++/for-11.C.jj	2013-06-21 08:38:19.000000000 +0200
+++ libgomp/testsuite/libgomp.c++/for-11.C	2013-09-04 15:28:21.542379101 +0200
@@ -75,10 +75,7 @@ int
 main ()
 {
   int err = 0;
-// FIXME: distribute construct must be closely nested
-// in teams region, but we don't handle target expansions
-// yet.  Enable when it works.
-// #pragma omp target teams reduction(|:err)
+  #pragma omp target teams reduction(|:err)
     {
       err |= test_d_normal ();
       err |= test_d_ds128_normal ();
--- libgomp/testsuite/libgomp.c++/target-1.C.jj	2013-09-04 18:55:02.294387946 +0200
+++ libgomp/testsuite/libgomp.c++/target-1.C	2013-09-04 18:54:56.430417616 +0200
@@ -0,0 +1 @@
+#include "../libgomp.c/target-1.c"
--- libgomp/testsuite/libgomp.c/for-3.c.jj	2013-07-06 19:44:29.000000000 +0200
+++ libgomp/testsuite/libgomp.c/for-3.c	2013-09-04 15:28:53.986223342 +0200
@@ -77,10 +77,7 @@ int
 main ()
 {
   int err = 0;
-/* FIXME: distribute construct must be closely nested
-   in teams region, but we don't handle target expansions
-   yet.  Enable when it works.  */
-/* #pragma omp target teams reduction(|:err) */
+  #pragma omp target teams reduction(|:err)
     {
       err |= test_d_normal ();
       err |= test_d_ds128_normal ();
--- libgomp/testsuite/libgomp.c/target-1.c.jj	2013-09-04 18:54:18.520609855 +0200
+++ libgomp/testsuite/libgomp.c/target-1.c	2013-09-04 18:54:10.000000000 +0200
@@ -0,0 +1,90 @@
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+fn1 (double *x, double *y, int z)
+{
+  int i;
+  for (i = 0; i < z; i++)
+    {
+      x[i] = i & 31;
+      y[i] = (i & 63) - 30;
+    }
+}
+
+#pragma omp declare target
+int tgtv = 6;
+int
+tgt (void)
+{
+  #pragma omp atomic update
+    tgtv++;
+  return 0;
+}
+#pragma omp end declare target
+
+double
+fn2 (int x, int y, int z)
+{
+  double b[1024], c[1024], s = 0;
+  int i, j;
+  fn1 (b, c, x);
+  #pragma omp target data map(to: b)
+  {
+    #pragma omp target map(tofrom: c)
+      #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
+	#pragma omp distribute dist_schedule(static, 4) collapse(1)
+	  for (j=0; j < x; j += y)
+	    #pragma omp parallel for reduction(+:s)
+	      for (i = j; i < j + y; i++)
+		tgt (), s += b[i] * c[i];
+    #pragma omp target update from(b, tgtv)
+  }
+  return s;
+}
+
+double
+fn3 (int x)
+{
+  double b[1024], c[1024], s = 0;
+  int i;
+  fn1 (b, c, x);
+  #pragma omp target map(to: b, c)
+    #pragma omp parallel for reduction(+:s)
+      for (i = 0; i < x; i++)
+	tgt (), s += b[i] * c[i];
+  return s;
+}
+
+double
+fn4 (int x, double *p)
+{
+  double b[1024], c[1024], d[1024], s = 0;
+  int i;
+  fn1 (b, c, x);
+  fn1 (d + x, p + x, x);
+  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
+    #pragma omp parallel for reduction(+:s)
+      for (i = 0; i < x; i++)
+	s += b[i] * c[i] + d[x + i] + p[x + i];
+  return s;
+}
+
+int
+main ()
+{
+  double a = fn2 (128, 4, 6);
+  int b = tgtv;
+  double c = fn3 (61);
+  #pragma omp target update from(tgtv)
+  int d = tgtv;
+  double e[1024];
+  double f = fn4 (64, e);
+  if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
+      || f != 8032.0)
+    abort ();
+  return 0;
+}
--- libgomp/target.c.jj	2013-09-04 14:58:30.488287438 +0200
+++ libgomp/target.c	2013-09-04 15:18:31.363404929 +0200
@@ -0,0 +1,83 @@
+/* Copyright (C) 2013 Free Software Foundation, Inc.
+   Contributed by Jakub Jelinek <jakub@redhat.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 maintainence of threads in response to team
+   creation and termination.  */
+
+#include "libgomp.h"
+#include <stdlib.h>
+#include <string.h>
+
+static int
+resolve_device (int device)
+{
+  return -1;
+}
+
+/* Called when encountering a target directive.  If DEVICE
+   is -1, it means use device-var ICV.  If it is -2 (or any other value
+   larger than last available hw device, use host fallback.
+   FN is address of host code, FNNAME corresponding name to lookup
+   in the target code.  HOSTADDRS, SIZES and KINDS are arrays
+   with MAPNUM entries, with addresses of the host objects,
+   sizes of the host objects (resp. for pointer kind pointer bias
+   and assumed sizeof (void *) size) and kinds.  */
+
+void
+GOMP_target (int device, void (*fn) (void *), const char *fnname,
+	     size_t mapnum, void **hostaddrs, size_t *sizes,
+	     unsigned char *kinds)
+{
+  if (resolve_device (device) == -1)
+    {
+      fn (hostaddrs);
+      return;
+    }
+}
+
+void
+GOMP_target_data (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
+		  unsigned char *kinds)
+{
+  if (resolve_device (device) == -1)
+    return;
+}
+
+void
+GOMP_target_end_data (void)
+{
+}
+
+void
+GOMP_target_update (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
+		    unsigned char *kinds)
+{
+  if (resolve_device (device) == -1)
+    return;
+}
+
+void
+GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
+{
+}
