diff mbox

[gomp4.1] Various accelerator updates from OpenMP 4.1

Message ID 20150729171907.GI1780@tucnak.redhat.com
State New
Headers show

Commit Message

Jakub Jelinek July 29, 2015, 5:19 p.m. UTC
On Fri, Jul 24, 2015 at 10:04:57PM +0200, Jakub Jelinek wrote:
> Another version.
> What to do with zero-length array sections vs. objects is still under heated
> debates, so target8.f90 keeps failing intermittently.

Here is a new version of the patch, with various additions (implemented
GOMP_MAP_FIRSTPRIVATE_INT I've talked about, it now handles use_device_ptr
and handles is_device_ptr with array decls (silly, but seems the accel folks
want it for some strange reason), etc.) and it special cases zero length
array sections rather than all zero length mappings.
The heated debates continue, so perhaps that part -
GOMP_MAP_ZERO_LEN_ARRAY_SECTION - will need reversion and replacement with
something else, we'll see.  This let's the testsuite pass for now except
for the two LTO ICEs, both without offloading (host fallback only) and with
Intel MIC offloading.  Committed to gomp-4_1-branch.

Ilya, I think now is the time to update your enter data/exit data patch.

2015-07-29  Jakub Jelinek  <jakub@redhat.com>

gcc/
	* tree.h (OMP_TARGET_COMBINED): Define.
	(OMP_CLAUSE_SET_MAP_KIND): Cast to unsigned int rather than unsigned
	char.
	(OMP_CLAUSE_MAP_PRIVATE,
	OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION): Define.
	* tree-core.h (struct tree_omp_clause): Change type of map_kind
	from unsigned char to unsigned int.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_0LEN_ARRAY.
	(enum omp_region_type): Add ORT_COMBINED_TARGET.
	(struct gimplify_omp_ctx): Add target_map_scalars_firstprivate,
	target_map_pointers_as_0len_arrays and
	target_firstprivatize_array_bases fields.
	(maybe_fold_stmt): Adjust check for ORT_TARGET for the addition of
	ORT_COMBINED_TARGET.
	(omp_notice_threadprivate_variable): Likewise.
	(omp_firstprivatize_variable): Likewise. 
	If ctx->target_map_scalars_firstprivate is set, firstprivatize
	as GOVD_FIRSTPRIVATE.
	(omp_add_variable): Allow map clause together with data sharing
	clauses.  For data sharing clause with VLA decl
	on omp target/target data don't add firstprivate for the pointer.
	(omp_notice_variable): Adjust check for ORT_TARGET for the addition
	of ORT_COMBINED_TARGET.  Handle implicit mapping of pointers
	as zero length array sections and
	ctx->target_map_scalars_firstprivate mapping of scalars as
	firstprivate data sharing.
	(gimplify_scan_omp_clauses): Initialize
	ctx->target_map_scalars_firstprivate,
	ctx->target_firstprivatize_array_bases and
	ctx->target_map_pointers_as_0len_arrays.  Add firstprivate for
	linear clause even to target region if combined.  Remove
	map clauses with GOMP_MAP_FIRSTPRIVATE_POINTER kind from
	OMP_TARGET_{,ENTER_,EXIT_}DATA.  For GOMP_MAP_FIRSTPRIVATE_POINTER
	map kind with non-INTEGER_CST OMP_CLAUSE_SIZE firstprivatize
	the bias.
	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_0LEN_ARRAY.
	If gimplify_omp_ctxp->target_firstprivatize_array_bases, use
	GOMP_MAP_FIRSTPRIVATE_POINTER map kind instead of
	GOMP_MAP_POINTER.
	(gimplify_adjust_omp_clauses): Adjust check for ORT_TARGET for the
	addition of ORT_COMBINED_TARGET.  Use
	GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER if
	ctx->target_firstprivatize_array_bases for VLAs.  Set
	OMP_CLAUSE_MAP_PRIVATE if both data sharing and map clause
	appear together.
	(gimplify_omp_workshare): Adjust check for ORT_TARGET for the
	addition of ORT_COMBINED_TARGET.  Use ORT_COMBINED_TARGET if
	OMP_TARGET_COMBINED.
	* omp-low.c (lookup_sfield): Change first argument to
	splay_tree_key, add overload with tree first argument.
	(maybe_lookup_field): Likewise.
	(build_sender_ref): Likewise.
	(scan_sharing_clauses): Handle VLAs in target firstprivate and
	is_device_ptr clauses.  Fix up variable shadowing.  Handle
	OMP_CLAUSE_USE_DEVICE_PTR.  Handle OMP_CLAUSE_MAP_PRIVATE.  Handle
	GOMP_MAP_FIRSTPRIVATE_POINTER map kind.
	(handle_simd_reference): Use get_name.
	(lower_rec_input_clauses): Likewise.  Use BUILT_IN_ALLOCA_WITH_ALIGN
	instead of BUILT_IN_ALLOCA.
	(lower_send_clauses): Use new lookup_sfield overload.
	(lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE_POINTER map kind.
	Handle OMP_CLAUSE_PRIVATE VLAs.  Handle OMP_CLAUSE_USE_DEVICE_PTR,
	handle arrays and references to arrays in OMP_CLAUSE_IS_DEVICE_PTR
	clause.  Handle OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION.
	* tree-pretty-print.c (dump_omp_clause): Handle
	GOMP_MAP_FIRSTPRIVATE_POINTER.
gcc/c/
	* c-tree.h (c_finish_omp_clauses): Add is_omp argument.
	* c-parser.c (c_parser_oacc_all_clauses, c_parser_omp_all_clauses,
	c_parser_oacc_cache, omp_split_clauses, c_parser_cilk_for): Adjust
	c_finish_omp_clauses callers.
	(c_parser_omp_target_data, c_parser_omp_target_enter_data,
	c_parser_omp_target_exit_data): Disallow GOMP_MAP_POINTER, allow
	GOMP_MAP_FIRSTPRIVATE_POINTER but don't set map_seen for it.
	(c_parser_omp_target): Set OMP_TARGET_COMBINED if combined.
	Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER.
	* c-typeck.c (handle_omp_array_sections): Add is_omp argument.
	Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if needed.
	Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER
	if is_omp.
	(c_finish_omp_clauses): Add is_omp argument, pass it down to
	handle_omp_array_sections.  Handle GOMP_MAP_FIRSTPRIVATE_POINTER.
	For is_device_ptr/use_device_ptr clauses allow ARRAY_TYPE.
gcc/cp/
	* parser.c (cp_parser_omp_target_data, cp_parser_omp_target_enter_data,
	cp_parser_omp_target_exit_data): Formatting fixes.  Disallow
	GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER but don't set
	map_seen for it.
	(cp_parser_omp_target): Set OMP_TARGET_COMBINED if combined.
	Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER.
	* semantics.c (handle_omp_array_sections): Add is_omp argument.
	Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if needed.
	Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER
	if is_omp.
	(finish_omp_clauses): Handle GOMP_MAP_FIRSTPRIVATE_POINTER.
	For is_device_ptr/use_device_ptr clauses allow ARRAY_TYPE
	and REFERENCE_TYPE to ARRAY_TYPE.
include/
	* gomp-constants.h (enum gomp_map_kind): Add
	GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_USE_DEVICE_PTR,
	GOMP_MAP_ZERO_LEN_ARRAY_SECTION and
	GOMP_MAP_FIRSTPRIVATE_POINTER.
libgomp/
	* libgomp.h (struct target_var_desc): Fix up comments
	about offset and length fields.
	* target.c (gomp_map_lookup): New function.
	(gomp_map_pointer): Use it.
	(gomp_map_vars): Handle GOMP_MAP_FIRSTPRIVATE_INT,
	GOMP_MAP_USE_DEVICE_PTR and GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
	Add tgt->list[i].offset for mappings with non-NULL
	tgt->list[i].key.
	(GOMP_target_41): Handle GOMP_MAP_FIRSTPRIVATE even
	for host fallback.
	(omp_target_is_present): Use gomp_map_lookup.
	(omp_target_associate_ptr): Likewise.
	(omp_target_disassociate_ptr): Likewise.
	* testsuite/libgomp.c++/target-2.C (fn2): Add map(tofrom: s).
	* testsuite/libgomp.c++/target-7.C: New test.
	* testsuite/libgomp.c++/target-8.C: New test.
	* testsuite/libgomp.c++/target-9.C: New test.
	* testsuite/libgomp.c/target-1.c (fn2, fn3, fn4): Add
	map(tofrom:s).
	* testsuite/libgomp.c/target-2.c (fn2, fn3, fn4): Likewise.
	* testsuite/libgomp.c/target-7.c (foo): Add map(h) where needed.
	* testsuite/libgomp.c/target-15.c: New test.
	* testsuite/libgomp.c/target-16.c: New test.
	* testsuite/libgomp.c/target-17.c: New test.
	* testsuite/libgomp.c/target-18.c: New test.
	* testsuite/libgomp.c/target-19.c: New test.
	* testsuite/libgomp.c/examples-4/e.51.3.c (gramSchmidt): Add
	map(tofrom:tmp).
	* testsuite/libgomp.c/examples-4/e.53.1.c (fib_wrapper): Add
	map(from:x).
	* testsuite/libgomp.c/examples-4/e.53.4.c (accum): Add
	map(tofrom:tmp).
	* testsuite/libgomp.c/examples-4/e.53.5.c (accum): Likewise.
	* testsuite/libgomp.c/examples-4/e.54.2.c (dotprod): Add
	map(tofrom: sum).
	* testsuite/libgomp.c/examples-4/e.54.3.c (dotprod): Likewise.
	* testsuite/libgomp.c/examples-4/e.54.4.c (dotprod): Likewise.
	* testsuite/libgomp.c/examples-4/e.57.1.c (main): Add
	map(from: c) and map(from: b, d) where needed.
	* testsuite/libgomp.c/examples-4/e.57.3.c (main): Add
	map(from: res).



	Jakub

Comments

Ilya Verbin Sept. 4, 2015, 6:07 p.m. UTC | #1
Hi!

It seems that there is a bug some here:

On Wed, Jul 29, 2015 at 19:19:07 +0200, Jakub Jelinek wrote:
> @@ -12918,6 +12989,28 @@ lower_omp_target (gimple_stmt_iterator *
>  	    var = var2;
>  	  }
>  
> +	if (offloaded
> +	    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
> +	  {
> +	    if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
> +	      {
> +		tree type = build_pointer_type (TREE_TYPE (var));
> +		tree new_var = lookup_decl (var, ctx);
> +		x = create_tmp_var_raw (type, get_name (new_var));
> +		gimple_add_tmp_var (x);
> +		x = build_simple_mem_ref (x);
> +		SET_DECL_VALUE_EXPR (new_var, x);
> +		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
> +	      }
> +	    continue;
> +	  }
> +
> +	if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
> +	  {
> +	    map_cnt++;
> +	    continue;
> +	  }
> +
>  	if (!maybe_lookup_field (var, ctx))
>  	  continue;
>  

Here is the reproducer:

#pragma omp declare target
int a[1];
#pragma omp end declare target

void foo ()
{
  #pragma omp target map(to: a[0:1])
    a;
}


lookup_decl (var, ctx) tries to lookup for 'a', but ctx->cb.decl_map->get ()
returns null-pointer.


$ gcc -fopenmp -c test.c

test.c: In function ‘foo’:
test.c:7:11: internal compiler error: Segmentation fault
   #pragma omp target map(to: a[0:1])
           ^
0xd27276 crash_signal
        gcc/toplev.c:352
0xbae3fa lookup_decl
        gcc/omp-low.c:1056
0xbe208c lower_omp_target
        gcc/omp-low.c:13362
0xbe8464 lower_omp_1
        gcc/omp-low.c:14504
0xbe8911 lower_omp
        gcc/omp-low.c:14592
0xbe8017 lower_omp_1
        gcc/omp-low.c:14436
0xbe8911 lower_omp
        gcc/omp-low.c:14592
0xbe808b lower_omp_1
        gcc/omp-low.c:14445
0xbe8911 lower_omp
        gcc/omp-low.c:14592
0xbe8ad9 execute_lower_omp
        gcc/omp-low.c:14630
0xbe8b7e execute
        gcc/omp-low.c:14667
Please submit a full bug report,
with preprocessed source if appropriate.
Please include the complete backtrace with any bug report.
See <http://gcc.gnu.org/bugs.html> for instructions.

  -- Ilya
Jakub Jelinek Sept. 4, 2015, 6:17 p.m. UTC | #2
On Fri, Sep 04, 2015 at 09:07:02PM +0300, Ilya Verbin wrote:
> It seems that there is a bug some here:

Thanks, will look at it on Monday.

> Here is the reproducer:
> 
> #pragma omp declare target
> int a[1];
> #pragma omp end declare target
> 
> void foo ()
> {
>   #pragma omp target map(to: a[0:1])
>     a;
> }

	Jakub
diff mbox

Patch

--- gcc/tree.h.jj	2015-07-16 17:56:41.000000000 +0200
+++ gcc/tree.h	2015-07-29 14:13:26.336307751 +0200
@@ -1341,6 +1341,11 @@  extern void protected_set_expr_location
 #define OMP_TEAMS_COMBINED(NODE) \
   (OMP_TEAMS_CHECK (NODE)->base.private_flag)
 
+/* True on an OMP_TARGET statement if it represents explicit
+   combined target teams, target parallel or target simd constructs.  */
+#define OMP_TARGET_COMBINED(NODE) \
+  (OMP_TARGET_CHECK (NODE)->base.private_flag)
+
 /* True if OMP_ATOMIC* is supposed to be sequentially consistent
    as opposed to relaxed.  */
 #define OMP_ATOMIC_SEQ_CST(NODE) \
@@ -1445,13 +1450,21 @@  extern void protected_set_expr_location
   ((enum gomp_map_kind) OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
 #define OMP_CLAUSE_SET_MAP_KIND(NODE, MAP_KIND) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
-   = (unsigned char) (MAP_KIND))
+   = (unsigned int) (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 corresponding
    OMP_CLAUSE_MAP with GOMP_MAP_POINTER 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)
+/* Nonzero if the same decl appears both in OMP_CLAUSE_MAP and either
+   OMP_CLAUSE_PRIVATE or OMP_CLAUSE_FIRSTPRIVATE.  */
+#define OMP_CLAUSE_MAP_PRIVATE(NODE) \
+  TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if this is a mapped array section, that might need special
+   treatment if OMP_CLAUSE_SIZE is zero.  */
+#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \
+  TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
 
 #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)
--- gcc/tree-core.h.jj	2015-07-17 09:30:44.000000000 +0200
+++ gcc/tree-core.h	2015-07-21 16:28:48.524156167 +0200
@@ -1354,7 +1354,7 @@  struct GTY(()) tree_omp_clause {
     enum omp_clause_schedule_kind  schedule_kind;
     enum omp_clause_depend_kind    depend_kind;
     /* See include/gomp-constants.h for enum gomp_map_kind's values.  */
-    unsigned char		   map_kind;
+    unsigned int		   map_kind;
     enum omp_clause_proc_bind_kind proc_bind_kind;
     enum tree_code                 reduction_code;
     enum omp_clause_linear_kind    linear_kind;
--- gcc/gimplify.c.jj	2015-07-16 17:56:41.000000000 +0200
+++ gcc/gimplify.c	2015-07-29 16:43:57.056823518 +0200
@@ -90,6 +90,8 @@  enum gimplify_omp_var_data
   /* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference.  */
   GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384,
 
+  GOVD_MAP_0LEN_ARRAY = 32768,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -110,6 +112,7 @@  enum omp_region_type
   ORT_TARGET_DATA = 16,
   /* Data region with offloading.  */
   ORT_TARGET = 32,
+  ORT_COMBINED_TARGET = 33,
   /* Dummy OpenMP region, used to disable expansion of
      DECL_VALUE_EXPRs in taskloop pre body.  */
   ORT_NONE = 64
@@ -156,6 +159,9 @@  struct gimplify_omp_ctx
   enum omp_region_type region_type;
   bool combined_loop;
   bool distribute;
+  bool target_map_scalars_firstprivate;
+  bool target_map_pointers_as_0len_arrays;
+  bool target_firstprivatize_array_bases;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -2260,7 +2266,7 @@  maybe_fold_stmt (gimple_stmt_iterator *g
 {
   struct gimplify_omp_ctx *ctx;
   for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
-    if (ctx->region_type == ORT_TARGET)
+    if ((ctx->region_type & ORT_TARGET) != 0)
       return false;
   return fold_stmt (gsi);
 }
@@ -5561,8 +5567,13 @@  omp_firstprivatize_variable (struct gimp
 	  else
 	    return;
 	}
-      else if (ctx->region_type == ORT_TARGET)
-	omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+      else if ((ctx->region_type & ORT_TARGET) != 0)
+	{
+	  if (ctx->target_map_scalars_firstprivate)
+	    omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
+	  else
+	    omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+	}
       else if (ctx->region_type != ORT_WORKSHARE
 	       && ctx->region_type != ORT_SIMD
 	       && ctx->region_type != ORT_TARGET_DATA)
@@ -5648,7 +5659,7 @@  omp_add_variable (struct gimplify_omp_ct
     flags |= GOVD_SEEN;
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (n != NULL && n->value != GOVD_ALIGNED)
+  if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
     {
       /* We shouldn't be re-adding the decl with the same data
 	 sharing class.  */
@@ -5678,6 +5689,9 @@  omp_add_variable (struct gimplify_omp_ct
 	    nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
 	  else if (flags & GOVD_PRIVATE)
 	    nflags = GOVD_PRIVATE;
+	  else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+		   && (flags & GOVD_FIRSTPRIVATE))
+	    nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
 	  else
 	    nflags = GOVD_FIRSTPRIVATE;
 	  nflags |= flags & GOVD_SEEN;
@@ -5746,7 +5760,7 @@  omp_notice_threadprivate_variable (struc
   struct gimplify_omp_ctx *octx;
 
   for (octx = ctx; octx; octx = octx->outer_context)
-    if (octx->region_type == ORT_TARGET)
+    if ((octx->region_type & ORT_TARGET) != 0)
       {
 	n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
 	if (n == NULL)
@@ -5810,19 +5824,66 @@  omp_notice_variable (struct gimplify_omp
     }
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (ctx->region_type == ORT_TARGET)
+  if ((ctx->region_type & ORT_TARGET) != 0)
     {
       ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
       if (n == NULL)
 	{
-	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+	  unsigned nflags = flags;
+	  if (ctx->target_map_pointers_as_0len_arrays
+	      || ctx->target_map_scalars_firstprivate)
+	    {
+	      bool is_declare_target = false;
+	      bool is_scalar = false;
+	      if (is_global_var (decl)
+		  && varpool_node::get_create (decl)->offloadable)
+		{
+		  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
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
+			break;
+		    }
+		  is_declare_target = octx == NULL;
+		}
+	      if (!is_declare_target && ctx->target_map_scalars_firstprivate)
+		{
+		  tree type = TREE_TYPE (decl);
+		  if (TREE_CODE (type) == REFERENCE_TYPE)
+		    type = TREE_TYPE (type);
+		  if (TREE_CODE (type) == COMPLEX_TYPE)
+		    type = TREE_TYPE (type);
+		  if (INTEGRAL_TYPE_P (type)
+		      || SCALAR_FLOAT_TYPE_P (type)
+		      || TREE_CODE (type) == POINTER_TYPE)
+		    is_scalar = true;
+		}
+	      if (is_declare_target)
+		;
+	      else if (ctx->target_map_pointers_as_0len_arrays
+		       && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+			   || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+			       && TREE_CODE (TREE_TYPE (TREE_TYPE (decl)))
+				  == POINTER_TYPE)))
+		nflags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
+	      else if (is_scalar)
+		nflags |= GOVD_FIRSTPRIVATE;
+	    }
+	  if (nflags == flags
+	      && !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);
+	      nflags |= GOVD_MAP | GOVD_EXPLICIT;
 	    }
-	  else
-	    omp_add_variable (ctx, decl, GOVD_MAP | flags);
+	  else if (nflags == flags)
+	    nflags |= GOVD_MAP;
+	  omp_add_variable (ctx, decl, nflags);
 	}
       else
 	{
@@ -6144,6 +6205,24 @@  gimplify_scan_omp_clauses (tree *list_p,
 
   ctx = new_omp_context (region_type);
   outer_ctx = ctx->outer_context;
+  if (code == OMP_TARGET && !lang_GNU_Fortran ())
+    {
+      ctx->target_map_pointers_as_0len_arrays = true;
+      /* FIXME: For Fortran we want to set this too, when
+	 the Fortran FE is updated to OpenMP 4.1.  */
+      ctx->target_map_scalars_firstprivate = true;
+    }
+  if (!lang_GNU_Fortran ())
+    switch (code)
+      {
+      case OMP_TARGET:
+      case OMP_TARGET_DATA:
+      case OMP_TARGET_ENTER_DATA:
+      case OMP_TARGET_EXIT_DATA:
+	ctx->target_firstprivatize_array_bases = true;
+      default:
+	break;
+      }
 
   while ((c = *list_p) != NULL)
     {
@@ -6290,11 +6369,18 @@  gimplify_scan_omp_clauses (tree *list_p,
 			   && ctx->region_type == ORT_WORKSHARE
 			   && octx == outer_ctx)
 		    flags = GOVD_SEEN | GOVD_SHARED;
+		  else if (octx
+			   && octx->region_type == ORT_COMBINED_TARGET)
+		    flags &= ~GOVD_LASTPRIVATE;
 		  else
 		    break;
-		  gcc_checking_assert (splay_tree_lookup (octx->variables,
-							  (splay_tree_key)
-							  decl) == NULL);
+		  splay_tree_node on
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  gcc_assert (on == NULL
+			      || (octx->region_type == ORT_COMBINED_TARGET
+				  && (on->value
+				      & GOVD_DATA_SHARE_CLASS) == 0));
 		  omp_add_variable (octx, decl, flags);
 		  if (octx->outer_context == NULL)
 		    break;
@@ -6319,10 +6405,24 @@  gimplify_scan_omp_clauses (tree *list_p,
 	case OMP_CLAUSE_MAP:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (error_operand_p (decl))
+	    remove = true;
+	  switch (code)
 	    {
-	      remove = true;
+	    case OMP_TARGET:
+	      break;
+	    case OMP_TARGET_DATA:
+	    case OMP_TARGET_ENTER_DATA:
+	    case OMP_TARGET_EXIT_DATA:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		/* For target {,enter ,exit }data only the array slice is
+		   mapped, but not the pointer to it.  */
+		remove = true;
+	      break;
+	    default:
 	      break;
 	    }
+	  if (remove)
+	    break;
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -6332,6 +6432,14 @@  gimplify_scan_omp_clauses (tree *list_p,
 	      remove = true;
 	      break;
 	    }
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+	    {
+	      OMP_CLAUSE_SIZE (c)
+		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL);
+	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	    }
 	  if (!DECL_P (decl))
 	    {
 	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
@@ -6643,7 +6751,10 @@  gimplify_scan_omp_clauses (tree *list_p,
 	case OMP_CLAUSE_NOGROUP:
 	case OMP_CLAUSE_THREADS:
 	case OMP_CLAUSE_SIMD:
+	  break;
+
 	case OMP_CLAUSE_DEFAULTMAP:
+	  ctx->target_map_scalars_firstprivate = false;
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -6759,6 +6870,30 @@  gimplify_adjust_omp_clauses_1 (splay_tre
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
     OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+  else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
+    {
+      tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_DECL (nc) = decl;
+      if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+	  && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE)
+	OMP_CLAUSE_DECL (clause)
+	  = build_simple_mem_ref_loc (input_location, decl);
+      OMP_CLAUSE_DECL (clause)
+	= build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause),
+		  build_int_cst (build_pointer_type (char_type_node), 0));
+      OMP_CLAUSE_SIZE (clause) = size_zero_node;
+      OMP_CLAUSE_SIZE (nc) = size_zero_node;
+      OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
+      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1;
+      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      OMP_CLAUSE_CHAIN (nc) = *list_p;
+      OMP_CLAUSE_CHAIN (clause) = nc;
+      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = ctx->outer_context;
+      gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
+		     pre_p, NULL, is_gimple_val, fb_rvalue);
+      gimplify_omp_ctxp = ctx;
+    }
   else if (code == OMP_CLAUSE_MAP)
     {
       OMP_CLAUSE_SET_MAP_KIND (clause,
@@ -6785,7 +6920,10 @@  gimplify_adjust_omp_clauses_1 (splay_tre
 				      OMP_CLAUSE_MAP);
 	  OMP_CLAUSE_DECL (nc) = decl;
 	  OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	  OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+	  if (gimplify_omp_ctxp->target_firstprivatize_array_bases)
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  else
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
 	  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
 	  OMP_CLAUSE_CHAIN (clause) = nc;
 	}
@@ -6910,12 +7048,14 @@  gimplify_adjust_omp_clauses (gimple_seq
 	  if (!DECL_P (decl))
 	    break;
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
-	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)
+	  if ((ctx->region_type & ORT_TARGET) != 0
+	      && !(n->value & GOVD_SEEN)
 	      && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
 	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
-		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
 	      /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
 		 for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -6935,17 +7075,33 @@  gimplify_adjust_omp_clauses (gimple_seq
 		  omp_notice_variable (ctx->outer_context,
 				       OMP_CLAUSE_SIZE (c), true);
 		}
-	      tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-					  OMP_CLAUSE_MAP);
-	      OMP_CLAUSE_DECL (nc) = decl;
-	      OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
-	      OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
-	      OMP_CLAUSE_CHAIN (c) = nc;
-	      c = nc;
+	      if (((ctx->region_type & ORT_TARGET) != 0
+		   || !ctx->target_firstprivatize_array_bases)
+		  && ((n->value & GOVD_SEEN) == 0
+		      || (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0))
+		{
+		  tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					      OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_DECL (nc) = decl;
+		  OMP_CLAUSE_SIZE (nc) = size_zero_node;
+		  if (ctx->target_firstprivatize_array_bases)
+		    OMP_CLAUSE_SET_MAP_KIND (nc,
+					     GOMP_MAP_FIRSTPRIVATE_POINTER);
+		  else
+		    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+		  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+		  OMP_CLAUSE_CHAIN (c) = nc;
+		  c = nc;
+		}
+	    }
+	  else
+	    {
+	      if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+		OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
+	      if ((n->value & GOVD_SEEN)
+		  && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
+		OMP_CLAUSE_MAP_PRIVATE (c) = 1;
 	    }
-	  else if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
-	    OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
 	  break;
 
 	case OMP_CLAUSE_TO:
@@ -7888,9 +8044,11 @@  gimplify_omp_workshare (tree *expr_p, gi
     case OMP_SINGLE:
       ort = ORT_WORKSHARE;
       break;
+    case OMP_TARGET:
+      ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
+      break;
     case OACC_KERNELS:
     case OACC_PARALLEL:
-    case OMP_TARGET:
       ort = ORT_TARGET;
       break;
     case OACC_DATA:
@@ -7905,7 +8063,7 @@  gimplify_omp_workshare (tree *expr_p, gi
     }
   gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
 			     TREE_CODE (expr));
-  if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+  if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
     {
       push_gimplify_context ();
       gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
--- gcc/omp-low.c.jj	2015-07-21 09:07:23.000000000 +0200
+++ gcc/omp-low.c	2015-07-29 16:13:33.209580272 +0200
@@ -1071,24 +1071,35 @@  lookup_field (tree var, omp_context *ctx
 }
 
 static inline tree
-lookup_sfield (tree var, omp_context *ctx)
+lookup_sfield (splay_tree_key key, omp_context *ctx)
 {
   splay_tree_node n;
   n = splay_tree_lookup (ctx->sfield_map
-			 ? ctx->sfield_map : ctx->field_map,
-			 (splay_tree_key) var);
+			 ? ctx->sfield_map : ctx->field_map, key);
   return (tree) n->value;
 }
 
 static inline tree
-maybe_lookup_field (tree var, omp_context *ctx)
+lookup_sfield (tree var, omp_context *ctx)
+{
+  return lookup_sfield ((splay_tree_key) var, ctx);
+}
+
+static inline tree
+maybe_lookup_field (splay_tree_key key, omp_context *ctx)
 {
   splay_tree_node n;
-  n = splay_tree_lookup (ctx->field_map, (splay_tree_key) var);
+  n = splay_tree_lookup (ctx->field_map, key);
   return n ? (tree) n->value : NULL_TREE;
 }
 
 static inline tree
+maybe_lookup_field (tree var, omp_context *ctx)
+{
+  return maybe_lookup_field ((splay_tree_key) var, ctx);
+}
+
+static inline tree
 lookup_oacc_reduction (const char *id, omp_context *ctx)
 {
   splay_tree_node n;
@@ -1359,12 +1370,18 @@  build_outer_var_ref (tree var, omp_conte
 /* Build tree nodes to access the field for VAR on the sender side.  */
 
 static tree
-build_sender_ref (tree var, omp_context *ctx)
+build_sender_ref (splay_tree_key key, omp_context *ctx)
 {
-  tree field = lookup_sfield (var, ctx);
+  tree field = lookup_sfield (key, ctx);
   return omp_build_component_ref (ctx->sender_decl, field);
 }
 
+static tree
+build_sender_ref (tree var, omp_context *ctx)
+{
+  return build_sender_ref ((splay_tree_key) var, ctx);
+}
+
 /* Add a new field for VAR inside the structure CTX->SENDER_DECL.  */
 
 static void
@@ -1908,6 +1925,17 @@  scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_LINEAR:
 	  decl = OMP_CLAUSE_DECL (c);
 	do_private:
+	  if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+	      && is_gimple_omp_offloaded (ctx->stmt))
+	    {
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+		install_var_field (decl, !is_reference (decl), 3, ctx);
+	      else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+		install_var_field (decl, true, 3, ctx);
+	      else
+		install_var_field (decl, false, 3, ctx);
+	    }
 	  if (is_variable_sized (decl))
 	    {
 	      if (is_task_ctx (ctx))
@@ -1930,10 +1958,6 @@  scan_sharing_clauses (tree clauses, omp_
 	      else if (!global)
 		install_var_field (decl, by_ref, 3, ctx);
 	    }
-	  else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
-		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
-		   && is_gimple_omp_offloaded (ctx->stmt))
-	    install_var_field (decl, !is_reference (decl), 3, ctx);
 	  install_var_local (decl, ctx);
 	  if (is_gimple_omp_oacc (ctx->stmt)
 	      && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
@@ -1944,9 +1968,9 @@  scan_sharing_clauses (tree clauses, omp_
 	      tree ptype = build_pointer_type (type);
 	      tree array = create_tmp_var (ptype,
 					   oacc_get_reduction_array_id (var));
-	      omp_context *c = (ctx->field_map ? ctx : ctx->outer);
-	      install_var_field (array, true, 3, c);
-	      install_var_local (array, c);
+	      omp_context *octx = (ctx->field_map ? ctx : ctx->outer);
+	      install_var_field (array, true, 3, octx);
+	      install_var_local (array, octx);
 
 	      /* Insert it into the current context.  */
 	      splay_tree_insert (ctx->reduction_map, (splay_tree_key)
@@ -1959,6 +1983,23 @@  scan_sharing_clauses (tree clauses, omp_
 	  break;
 
 	case OMP_CLAUSE_USE_DEVICE_PTR:
+	  decl = OMP_CLAUSE_DECL (c);
+	  if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+	    install_var_field (decl, true, 3, ctx);
+	  else
+	    install_var_field (decl, false, 3, ctx);
+	  if (DECL_SIZE (decl)
+	      && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+	    {
+	      tree decl2 = DECL_VALUE_EXPR (decl);
+	      gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+	      decl2 = TREE_OPERAND (decl2, 0);
+	      gcc_assert (DECL_P (decl2));
+	      install_var_local (decl2, ctx);
+	    }
+	  install_var_local (decl, ctx);
+	  break;
+
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  goto do_private;
@@ -2025,6 +2066,21 @@  scan_sharing_clauses (tree clauses, omp_
 		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	    {
+	      if (DECL_SIZE (decl)
+		  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+		{
+		  tree decl2 = DECL_VALUE_EXPR (decl);
+		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+		  decl2 = TREE_OPERAND (decl2, 0);
+		  gcc_assert (DECL_P (decl2));
+		  install_var_local (decl2, ctx);
+		}
+	      install_var_local (decl, ctx);
+	      break;
+	    }
 	  if (DECL_P (decl))
 	    {
 	      if (DECL_SIZE (decl)
@@ -2034,7 +2090,11 @@  scan_sharing_clauses (tree clauses, omp_
 		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
 		  decl2 = TREE_OPERAND (decl2, 0);
 		  gcc_assert (DECL_P (decl2));
-		  install_var_field (decl2, true, 3, ctx);
+		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		      && OMP_CLAUSE_MAP_PRIVATE (c))
+		    install_var_field (decl2, true, 11, ctx);
+		  else
+		    install_var_field (decl2, true, 3, ctx);
 		  install_var_local (decl2, ctx);
 		  install_var_local (decl, ctx);
 		}
@@ -2045,6 +2105,9 @@  scan_sharing_clauses (tree clauses, omp_
 		      && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 		    install_var_field (decl, true, 7, ctx);
+		  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			   && OMP_CLAUSE_MAP_PRIVATE (c))
+		    install_var_field (decl, true, 11, ctx);
 		  else
 		    install_var_field (decl, true, 3, ctx);
 		  if (is_gimple_omp_offloaded (ctx->stmt))
@@ -2147,11 +2210,23 @@  scan_sharing_clauses (tree clauses, omp_
 	  /* FALLTHRU */
 	case OMP_CLAUSE_PRIVATE:
 	case OMP_CLAUSE_LINEAR:
-	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_variable_sized (decl))
-	    install_var_local (decl, ctx);
+	    {
+	      if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+		   || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+		  && is_gimple_omp_offloaded (ctx->stmt))
+		{
+		  tree decl2 = DECL_VALUE_EXPR (decl);
+		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+		  decl2 = TREE_OPERAND (decl2, 0);
+		  gcc_assert (DECL_P (decl2));
+		  install_var_local (decl2, ctx);
+		  fixup_remapped_decl (decl2, ctx, false);
+		}
+	      install_var_local (decl, ctx);
+	    }
 	  fixup_remapped_decl (decl, ctx,
 			       OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
 			       && OMP_CLAUSE_PRIVATE_DEBUG (c));
@@ -2201,7 +2276,8 @@  scan_sharing_clauses (tree clauses, omp_
 	    break;
 	  if (DECL_P (decl))
 	    {
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		   || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 		  && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
 		  && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
 		{
@@ -2255,6 +2331,7 @@  scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_NOGROUP:
 	case OMP_CLAUSE_DEFAULTMAP:
+	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
@@ -3924,11 +4001,8 @@  handle_simd_reference (location_t loc, t
   tree z = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard)));
   if (TREE_CONSTANT (z))
     {
-      const char *name = NULL;
-      if (DECL_NAME (new_vard))
-	name = IDENTIFIER_POINTER (DECL_NAME (new_vard));
-
-      z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), name);
+      z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)),
+			      get_name (new_vard));
       gimple_add_tmp_var (z);
       TREE_ADDRESSABLE (z) = 1;
       z = build_fold_addr_expr_loc (loc, z);
@@ -4127,9 +4201,7 @@  lower_rec_input_clauses (tree clauses, g
 	      tree type = TREE_TYPE (d);
 	      gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
 	      tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
-	      const char *name = NULL;
-	      if (DECL_NAME (orig_var))
-		name = IDENTIFIER_POINTER (DECL_NAME (orig_var));
+	      const char *name = get_name (orig_var);
 	      if (TREE_CONSTANT (v))
 		{
 		  x = create_tmp_var_raw (type, name);
@@ -4139,7 +4211,8 @@  lower_rec_input_clauses (tree clauses, g
 		}
 	      else
 		{
-		  tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
+		  tree atmp
+		    = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
 		  tree t = maybe_lookup_decl (v, ctx);
 		  if (t)
 		    v = t;
@@ -4152,7 +4225,8 @@  lower_rec_input_clauses (tree clauses, g
 		  t = fold_build2_loc (clause_loc, MULT_EXPR,
 				       TREE_TYPE (v), t,
 				       TYPE_SIZE_UNIT (TREE_TYPE (type)));
-		  x = build_call_expr_loc (clause_loc, atmp, 1, t);
+		  tree al = size_int (TYPE_ALIGN (TREE_TYPE (type)));
+		  x = build_call_expr_loc (clause_loc, atmp, 2, t, al);
 		}
 
 	      tree ptype = build_pointer_type (TREE_TYPE (type));
@@ -4362,8 +4436,9 @@  lower_rec_input_clauses (tree clauses, g
 		  x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
 
 		  /* void *tmp = __builtin_alloca */
-		  atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		  stmt = gimple_build_call (atmp, 1, x);
+		  atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		  stmt = gimple_build_call (atmp, 2, x,
+					    size_int (DECL_ALIGN (var)));
 		  tmp = create_tmp_var_raw (ptr_type_node);
 		  gimple_add_tmp_var (tmp);
 		  gimple_call_set_lhs (stmt, tmp);
@@ -4400,12 +4475,8 @@  lower_rec_input_clauses (tree clauses, g
 		    x = NULL_TREE;
 		  else
 		    {
-		      const char *name = NULL;
-		      if (DECL_NAME (var))
-			name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
 		      x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
-					      name);
+					      get_name (var));
 		      gimple_add_tmp_var (x);
 		      TREE_ADDRESSABLE (x) = 1;
 		      x = build_fold_addr_expr_loc (clause_loc, x);
@@ -4413,8 +4484,11 @@  lower_rec_input_clauses (tree clauses, g
 		}
 	      else
 		{
-		  tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		  x = build_call_expr_loc (clause_loc, atmp, 1, x);
+		  tree atmp
+		    = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		  tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+		  tree al = size_int (TYPE_ALIGN (rtype));
+		  x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
 		}
 
 	      if (x)
@@ -5489,11 +5563,7 @@  lower_send_clauses (tree clauses, gimple
 	  /* Handle taskloop firstprivate/lastprivate, where the
 	     lastprivate on GIMPLE_OMP_TASK is represented as
 	     OMP_CLAUSE_SHARED_FIRSTPRIVATE.  */
-	  tree f
-	    = (tree)
-	      splay_tree_lookup (ctx->sfield_map
-				 ? ctx->sfield_map : ctx->field_map,
-				 (splay_tree_key) &DECL_UID (val))->value;
+	  tree f = lookup_sfield ((splay_tree_key) &DECL_UID (val), ctx);
 	  x = omp_build_component_ref (ctx->sender_decl, f);
 	  if (use_pointer_for_field (val, ctx))
 	    var = build_fold_addr_expr (var);
@@ -12883,6 +12953,7 @@  lower_omp_target (gimple_stmt_iterator *
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
@@ -12918,6 +12989,28 @@  lower_omp_target (gimple_stmt_iterator *
 	    var = var2;
 	  }
 
+	if (offloaded
+	    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	  {
+	    if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+	      {
+		tree type = build_pointer_type (TREE_TYPE (var));
+		tree new_var = lookup_decl (var, ctx);
+		x = create_tmp_var_raw (type, get_name (new_var));
+		gimple_add_tmp_var (x);
+		x = build_simple_mem_ref (x);
+		SET_DECL_VALUE_EXPR (new_var, x);
+		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	      }
+	    continue;
+	  }
+
+	if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
+	  {
+	    map_cnt++;
+	    continue;
+	  }
+
 	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
@@ -12925,6 +13018,7 @@  lower_omp_target (gimple_stmt_iterator *
 	  {
 	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+
 	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
@@ -12936,14 +13030,70 @@  lower_omp_target (gimple_stmt_iterator *
 	break;
 
       case OMP_CLAUSE_FIRSTPRIVATE:
-      case OMP_CLAUSE_IS_DEVICE_PTR:
 	map_cnt++;
 	var = OMP_CLAUSE_DECL (c);
 	if (!is_reference (var)
 	    && !is_gimple_reg_type (TREE_TYPE (var)))
 	  {
-	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+	    if (is_variable_sized (var))
+	      {
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_pvar = lookup_decl (pvar, ctx);
+		x = build_fold_indirect_ref (new_pvar);
+		TREE_THIS_NOTRAP (x) = 1;
+	      }
+	    else
+	      x = build_receiver_ref (var, true, ctx);
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	break;
+
+      case OMP_CLAUSE_PRIVATE:
+	var = OMP_CLAUSE_DECL (c);
+	if (is_variable_sized (var))
+	  {
+	    tree new_var = lookup_decl (var, ctx);
+	    tree pvar = DECL_VALUE_EXPR (var);
+	    gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+	    pvar = TREE_OPERAND (pvar, 0);
+	    gcc_assert (DECL_P (pvar));
+	    tree new_pvar = lookup_decl (pvar, ctx);
+	    x = build_fold_indirect_ref (new_pvar);
+	    TREE_THIS_NOTRAP (x) = 1;
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	break;
+
+      case OMP_CLAUSE_USE_DEVICE_PTR:
+      case OMP_CLAUSE_IS_DEVICE_PTR:
+	var = OMP_CLAUSE_DECL (c);
+	map_cnt++;
+	if (is_variable_sized (var))
+	  {
+	    tree new_var = lookup_decl (var, ctx);
+	    tree pvar = DECL_VALUE_EXPR (var);
+	    gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+	    pvar = TREE_OPERAND (pvar, 0);
+	    gcc_assert (DECL_P (pvar));
+	    tree new_pvar = lookup_decl (pvar, ctx);
+	    x = build_fold_indirect_ref (new_pvar);
+	    TREE_THIS_NOTRAP (x) = 1;
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+	  {
+	    tree new_var = lookup_decl (var, ctx);
+	    tree type = build_pointer_type (TREE_TYPE (var));
+	    x = create_tmp_var_raw (type, get_name (new_var));
+	    gimple_add_tmp_var (x);
+	    x = build_simple_mem_ref (x);
 	    SET_DECL_VALUE_EXPR (new_var, x);
 	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
 	  }
@@ -13013,7 +13163,7 @@  lower_omp_target (gimple_stmt_iterator *
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
-	    tree ovar, nc, s, purpose, var, x;
+	    tree ovar, nc, s, purpose, var, x, type;
 	    unsigned int talign;
 
 	  default:
@@ -13044,6 +13194,10 @@  lower_omp_target (gimple_stmt_iterator *
 	      }
 	    else
 	      {
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_KIND (c)
+		       == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		  break;
 		if (DECL_SIZE (ovar)
 		    && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
 		  {
@@ -13053,7 +13207,14 @@  lower_omp_target (gimple_stmt_iterator *
 		    gcc_assert (DECL_P (ovar2));
 		    ovar = ovar2;
 		  }
-		if (!maybe_lookup_field (ovar, ctx))
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_PRIVATE (c))
+		  {
+		    if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar),
+					     ctx))
+		      continue;
+		  }
+		else if (!maybe_lookup_field (ovar, ctx))
 		  continue;
 	      }
 
@@ -13063,7 +13224,12 @@  lower_omp_target (gimple_stmt_iterator *
 	    if (nc)
 	      {
 		var = lookup_decl_in_outer_ctx (ovar, ctx);
-		x = build_sender_ref (ovar, ctx);
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_PRIVATE (c))
+		  x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar),
+					ctx);
+		else
+		  x = build_sender_ref (ovar, ctx);
 		if (maybe_lookup_oacc_reduction (var, ctx))
 		  {
 		    gcc_checking_assert (offloaded
@@ -13101,7 +13267,7 @@  lower_omp_target (gimple_stmt_iterator *
 			 || map_kind == GOMP_MAP_FORCE_DEVICEPTR)
 			&& !TYPE_READONLY (TREE_TYPE (var)))
 		      {
-			x = build_sender_ref (ovar, ctx);
+			x = unshare_expr (x);
 			x = build_simple_mem_ref (x);
 			gimplify_assign (var, x, &olist);
 		      }
@@ -13121,35 +13287,74 @@  lower_omp_target (gimple_stmt_iterator *
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    unsigned HOST_WIDE_INT tkind;
+	    unsigned HOST_WIDE_INT tkind, tkind_zero;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_MAP:
 		tkind = OMP_CLAUSE_MAP_KIND (c);
+		tkind_zero = tkind;
+		if (OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c))
+		  switch (tkind)
+		    {
+		    case GOMP_MAP_ALLOC:
+		    case GOMP_MAP_TO:
+		    case GOMP_MAP_FROM:
+		    case GOMP_MAP_TOFROM:
+		    case GOMP_MAP_ALWAYS_TO:
+		    case GOMP_MAP_ALWAYS_FROM:
+		    case GOMP_MAP_ALWAYS_TOFROM:
+		      tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
+		      break;
+		    default:
+		      break;
+		    }
+		if (tkind_zero != tkind)
+		  {
+		    if (integer_zerop (s))
+		      tkind = tkind_zero;
+		    else if (integer_nonzerop (s))
+		      tkind_zero = tkind;
+		  }
 		break;
 	      case OMP_CLAUSE_TO:
 		tkind = GOMP_MAP_TO;
+		tkind_zero = tkind;
 		break;
 	      case OMP_CLAUSE_FROM:
 		tkind = GOMP_MAP_FROM;
+		tkind_zero = tkind;
 		break;
 	      default:
 		gcc_unreachable ();
 	      }
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
+	    gcc_checking_assert (tkind_zero
+				 < (HOST_WIDE_INT_C (1U) << talign_shift));
 	    talign = ceil_log2 (talign);
 	    tkind |= talign << talign_shift;
+	    tkind_zero |= talign << talign_shift;
 	    gcc_checking_assert (tkind
 				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
-	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
-				    build_int_cstu (tkind_type, tkind));
+	    gcc_checking_assert (tkind_zero
+				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+	    if (tkind == tkind_zero)
+	      x = build_int_cstu (tkind_type, tkind);
+	    else
+	      {
+		TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0;
+		x = build3 (COND_EXPR, tkind_type,
+			    fold_build2 (EQ_EXPR, boolean_type_node,
+					 unshare_expr (s), size_zero_node),
+			    build_int_cstu (tkind_type, tkind_zero),
+			    build_int_cstu (tkind_type, tkind));
+	      }
+	    CONSTRUCTOR_APPEND_ELT (vkind, purpose, x);
 	    if (nc && nc != c)
 	      c = nc;
 	    break;
 
 	  case OMP_CLAUSE_FIRSTPRIVATE:
-	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    ovar = OMP_CLAUSE_DECL (c);
 	    if (is_reference (ovar))
 	      talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@@ -13157,7 +13362,24 @@  lower_omp_target (gimple_stmt_iterator *
 	      talign = DECL_ALIGN_UNIT (ovar);
 	    var = lookup_decl_in_outer_ctx (ovar, ctx);
 	    x = build_sender_ref (ovar, ctx);
-	    if (is_reference (var))
+	    tkind = GOMP_MAP_FIRSTPRIVATE;
+	    type = TREE_TYPE (ovar);
+	    if (is_reference (ovar))
+	      type = TREE_TYPE (type);
+	    if ((INTEGRAL_TYPE_P (type)
+		  && TYPE_PRECISION (type) <= POINTER_SIZE)
+		|| TREE_CODE (type) == POINTER_TYPE)
+	      {
+		tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+		tree t = var;
+		if (is_reference (var))
+		  t = build_simple_mem_ref (var);
+		if (TREE_CODE (type) != POINTER_TYPE)
+		  t = fold_convert (pointer_sized_int_node, t);
+		t = fold_convert (TREE_TYPE (x), t);
+		gimplify_assign (x, t, &ilist);
+	      }
+	    else if (is_reference (var))
 	      gimplify_assign (x, var, &ilist);
 	    else if (is_gimple_reg (var))
 	      {
@@ -13172,7 +13394,9 @@  lower_omp_target (gimple_stmt_iterator *
 		var = build_fold_addr_expr (var);
 		gimplify_assign (x, var, &ilist);
 	      }
-	    if (is_reference (var))
+	    if (tkind == GOMP_MAP_FIRSTPRIVATE_INT)
+	      s = size_int (0);
+	    else if (is_reference (var))
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
 	    else
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
@@ -13182,7 +13406,6 @@  lower_omp_target (gimple_stmt_iterator *
 	    if (TREE_CODE (s) != INTEGER_CST)
 	      TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
 
-	    tkind = GOMP_MAP_FIRSTPRIVATE;
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
 	    talign = ceil_log2 (talign);
@@ -13192,6 +13415,40 @@  lower_omp_target (gimple_stmt_iterator *
 	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
 				    build_int_cstu (tkind_type, tkind));
 	    break;
+
+	  case OMP_CLAUSE_USE_DEVICE_PTR:
+	  case OMP_CLAUSE_IS_DEVICE_PTR:
+	    ovar = OMP_CLAUSE_DECL (c);
+	    var = lookup_decl_in_outer_ctx (ovar, ctx);
+	    x = build_sender_ref (ovar, ctx);
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+	      tkind = GOMP_MAP_USE_DEVICE_PTR;
+	    else
+	      tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+	    type = TREE_TYPE (ovar);
+	    if (TREE_CODE (type) == ARRAY_TYPE)
+	      var = build_fold_addr_expr (var);
+	    else
+	      {
+		if (is_reference (ovar))
+		  {
+		    type = TREE_TYPE (type);
+		    if (TREE_CODE (type) != ARRAY_TYPE)
+		      var = build_simple_mem_ref (var);
+		    var = fold_convert (TREE_TYPE (x), var);
+		  }
+	      }
+	    gimplify_assign (x, var, &ilist);
+	    s = size_int (0);
+	    purpose = size_int (map_idx++);
+	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+	    gcc_checking_assert (tkind
+				 < (HOST_WIDE_INT_C (1U) << talign_shift));
+	    gcc_checking_assert (tkind
+				 <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+	    CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+				    build_int_cstu (tkind_type, tkind));
+	    break;
 	  }
 
       gcc_assert (map_idx == map_cnt);
@@ -13200,21 +13457,22 @@  lower_omp_target (gimple_stmt_iterator *
 	= 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 (TREE_TYPE (TREE_VEC_ELT (t, 1)),
-					    NULL);
-	  TREE_THIS_VOLATILE (clobber) = 1;
-	  gimple_seq_add_stmt (&olist,
-			       gimple_build_assign (TREE_VEC_ELT (t, 1),
-						    clobber));
-	}
+      for (int i = 1; i <= 2; i++)
+	if (!TREE_STATIC (TREE_VEC_ELT (t, i)))
+	  {
+	    gimple_seq initlist = NULL;
+	    force_gimple_operand (build1 (DECL_EXPR, void_type_node,
+					  TREE_VEC_ELT (t, i)),
+				  &initlist, true, NULL_TREE);
+	    gimple_seq_add_seq (&ilist, initlist);
+
+	    tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, i)),
+					      NULL);
+	    TREE_THIS_VOLATILE (clobber) = 1;
+	    gimple_seq_add_stmt (&olist,
+				 gimple_build_assign (TREE_VEC_ELT (t, i),
+						      clobber));
+	  }
 
       tree clobber = build_constructor (ctx->record_type, NULL);
       TREE_THIS_VOLATILE (clobber) = 1;
@@ -13237,22 +13495,64 @@  lower_omp_target (gimple_stmt_iterator *
 	  		   gimple_build_assign (ctx->receiver_decl, t));
     }
 
-  if (offloaded)
+  if (offloaded || data_region)
     {
+      tree prev = NULL_TREE;
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
-	    tree var;
+	    tree var, x;
 	  default:
 	    break;
 	  case OMP_CLAUSE_FIRSTPRIVATE:
-	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    var = OMP_CLAUSE_DECL (c);
 	    if (is_reference (var)
 		|| is_gimple_reg_type (TREE_TYPE (var)))
 	      {
 		tree new_var = lookup_decl (var, ctx);
-		tree x = build_receiver_ref (var, !is_reference (var), ctx);
+		tree type;
+		type = TREE_TYPE (var);
+		if (is_reference (var))
+		  type = TREE_TYPE (type);
+		if ((INTEGRAL_TYPE_P (type)
+		     && TYPE_PRECISION (type) <= POINTER_SIZE)
+		    || TREE_CODE (type) == POINTER_TYPE)
+		  {
+		    x = build_receiver_ref (var, false, ctx);
+		    if (TREE_CODE (type) != POINTER_TYPE)
+		      x = fold_convert (pointer_sized_int_node, x);
+		    x = fold_convert (type, x);
+		    gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+				   fb_rvalue);
+		    if (is_reference (var))
+		      {
+			tree v = create_tmp_var_raw (type, get_name (var));
+			gimple_add_tmp_var (v);
+			TREE_ADDRESSABLE (v) = 1;
+			gimple_seq_add_stmt (&new_body,
+					     gimple_build_assign (v, x));
+			x = build_fold_addr_expr (v);
+		      }
+		    gimple_seq_add_stmt (&new_body,
+					 gimple_build_assign (new_var, x));
+		  }
+		else
+		  {
+		    x = build_receiver_ref (var, !is_reference (var), ctx);
+		    gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+				   fb_rvalue);
+		    gimple_seq_add_stmt (&new_body,
+					 gimple_build_assign (new_var, x));
+		  }
+	      }
+	    else if (is_variable_sized (var))
+	      {
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_var = lookup_decl (pvar, ctx);
+		x = build_receiver_ref (var, false, ctx);
 		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
 		gimple_seq_add_stmt (&new_body,
 				     gimple_build_assign (new_var, x));
@@ -13264,23 +13564,22 @@  lower_omp_target (gimple_stmt_iterator *
 	      {
 		location_t clause_loc = OMP_CLAUSE_LOCATION (c);
 		tree new_var = lookup_decl (var, ctx);
-		tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
+		x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
 		if (TREE_CONSTANT (x))
 		  {
-		    const char *name = NULL;
-		    if (DECL_NAME (var))
-		      name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
 		    x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
-					    name);
+					    get_name (var));
 		    gimple_add_tmp_var (x);
 		    TREE_ADDRESSABLE (x) = 1;
 		    x = build_fold_addr_expr_loc (clause_loc, x);
 		  }
 		else
 		  {
-		    tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		    x = build_call_expr_loc (clause_loc, atmp, 1, x);
+		    tree atmp
+		      = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		    tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+		    tree al = size_int (TYPE_ALIGN (rtype));
+		    x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
 		  }
 
 		x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
@@ -13289,9 +13588,169 @@  lower_omp_target (gimple_stmt_iterator *
 				     gimple_build_assign (new_var, x));
 	      }
 	    break;
+	  case OMP_CLAUSE_USE_DEVICE_PTR:
+	  case OMP_CLAUSE_IS_DEVICE_PTR:
+	    var = OMP_CLAUSE_DECL (c);
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+	      x = build_sender_ref (var, ctx);
+	    else
+	      x = build_receiver_ref (var, false, ctx);
+	    if (is_variable_sized (var))
+	      {
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_var = lookup_decl (pvar, ctx);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+	      }
+	    else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+	      {
+		tree new_var = lookup_decl (var, ctx);
+		new_var = DECL_VALUE_EXPR (new_var);
+		gcc_assert (TREE_CODE (new_var) == MEM_REF);
+		new_var = TREE_OPERAND (new_var, 0);
+		gcc_assert (DECL_P (new_var));
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+	      }
+	    else
+	      {
+		tree type = TREE_TYPE (var);
+		tree new_var = lookup_decl (var, ctx);
+		if (is_reference (var))
+		  {
+		    type = TREE_TYPE (type);
+		    if (TREE_CODE (type) != ARRAY_TYPE)
+		      {
+			tree v = create_tmp_var_raw (type, get_name (var));
+			gimple_add_tmp_var (v);
+			TREE_ADDRESSABLE (v) = 1;
+			x = fold_convert (type, x);
+			gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+				       fb_rvalue);
+			gimple_seq_add_stmt (&new_body,
+					     gimple_build_assign (v, x));
+			x = build_fold_addr_expr (v);
+		      }
+		  }
+		x = fold_convert (TREE_TYPE (new_var), x);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+	      }
+	    break;
+	  }
+      /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass,
+	 so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
+	 are already handled.  */
+      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	    tree var;
+	  default:
+	    break;
+	  case OMP_CLAUSE_MAP:
+	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	      {
+		location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+		gcc_assert (prev);
+		var = OMP_CLAUSE_DECL (c);
+		if (DECL_SIZE (var)
+		    && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+		  {
+		    tree var2 = DECL_VALUE_EXPR (var);
+		    gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+		    var2 = TREE_OPERAND (var2, 0);
+		    gcc_assert (DECL_P (var2));
+		    var = var2;
+		  }
+		tree new_var = lookup_decl (var, ctx), x;
+		tree type = TREE_TYPE (new_var);
+		bool is_ref = is_reference (var);
+		bool ref_to_array = false;
+		if (is_ref)
+		  {
+		    type = TREE_TYPE (type);
+		    if (TREE_CODE (type) == ARRAY_TYPE)
+		      {
+			type = build_pointer_type (type);
+			ref_to_array = true;
+		      }
+		  }
+		else if (TREE_CODE (type) == ARRAY_TYPE)
+		  {
+		    tree decl2 = DECL_VALUE_EXPR (new_var);
+		    gcc_assert (TREE_CODE (decl2) == MEM_REF);
+		    decl2 = TREE_OPERAND (decl2, 0);
+		    gcc_assert (DECL_P (decl2));
+		    new_var = decl2;
+		    type = TREE_TYPE (new_var);
+		  }
+		x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx);
+		x = fold_convert_loc (clause_loc, type, x);
+		if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
+		  {
+		    tree bias = OMP_CLAUSE_SIZE (c);
+		    if (DECL_P (bias))
+		      bias = lookup_decl (bias, ctx);
+		    bias = fold_convert_loc (clause_loc, sizetype, bias);
+		    bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype,
+					    bias);
+		    x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+					 TREE_TYPE (x), x, bias);
+		  }
+		if (ref_to_array)
+		  x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		if (is_ref && !ref_to_array)
+		  {
+		    tree t = create_tmp_var_raw (type, get_name (var));
+		    gimple_add_tmp_var (t);
+		    TREE_ADDRESSABLE (t) = 1;
+		    gimple_seq_add_stmt (&new_body,
+					 gimple_build_assign (t, x));
+		    x = build_fold_addr_expr_loc (clause_loc, t);
+		  }
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+		prev = NULL_TREE;
+	      }
+	    else if (OMP_CLAUSE_CHAIN (c)
+		     && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+			== OMP_CLAUSE_MAP
+		     && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			== GOMP_MAP_FIRSTPRIVATE_POINTER)
+	      prev = c;
+	    break;
+	  case OMP_CLAUSE_PRIVATE:
+	    var = OMP_CLAUSE_DECL (c);
+	    if (is_variable_sized (var))
+	      {
+		location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+		tree new_var = lookup_decl (var, ctx);
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_pvar = lookup_decl (pvar, ctx);
+		tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		tree al = size_int (DECL_ALIGN (var));
+		tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
+		x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+		x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_pvar, x));
+	      }
+	    break;
 	  }
       gimple_seq_add_seq (&new_body, tgt_body);
-      new_body = maybe_catch_exception (new_body);
+      if (offloaded)
+	new_body = maybe_catch_exception (new_body);
     }
   else if (data_region)
     new_body = tgt_body;
--- gcc/tree-pretty-print.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/tree-pretty-print.c	2015-07-22 13:53:51.406065024 +0200
@@ -639,6 +639,9 @@  dump_omp_clause (pretty_printer *pp, tre
 	case GOMP_MAP_RELEASE:
 	  pp_string (pp, "release");
 	  break;
+	case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	  pp_string (pp, "firstprivate");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -649,7 +652,9 @@  dump_omp_clause (pretty_printer *pp, tre
       if (OMP_CLAUSE_SIZE (clause))
 	{
 	  if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
-	      && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER)
+	      && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER
+		  || OMP_CLAUSE_MAP_KIND (clause)
+		     == GOMP_MAP_FIRSTPRIVATE_POINTER))
 	    pp_string (pp, " [pointer assign, bias: ");
 	  else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)
--- gcc/c/c-tree.h.jj	2015-07-01 12:50:49.000000000 +0200
+++ gcc/c/c-tree.h	2015-07-22 12:47:49.185826677 +0200
@@ -649,7 +649,7 @@  extern tree c_begin_omp_task (void);
 extern tree c_finish_omp_task (location_t, tree, tree);
 extern void c_finish_omp_cancel (location_t, tree);
 extern void c_finish_omp_cancellation_point (location_t, tree);
-extern tree c_finish_omp_clauses (tree, bool = false);
+extern tree c_finish_omp_clauses (tree, bool, bool = false);
 extern tree c_build_va_arg (location_t, tree, tree);
 extern tree c_finish_transaction (location_t, tree, int);
 extern bool c_tree_equal (tree, tree);
--- gcc/c/c-parser.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/c/c-parser.c	2015-07-23 12:51:02.000000000 +0200
@@ -12435,7 +12435,7 @@  c_parser_oacc_all_clauses (c_parser *par
   c_parser_skip_to_pragma_eol (parser);
 
   if (finish_p)
-    return c_finish_omp_clauses (clauses);
+    return c_finish_omp_clauses (clauses, false);
 
   return clauses;
 }
@@ -12720,8 +12720,8 @@  c_parser_omp_all_clauses (c_parser *pars
   if (finish_p)
     {
       if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
-	return c_finish_omp_clauses (clauses, true);
-      return c_finish_omp_clauses (clauses);
+	return c_finish_omp_clauses (clauses, true, true);
+      return c_finish_omp_clauses (clauses, true);
     }
 
   return clauses;
@@ -12755,7 +12755,7 @@  c_parser_oacc_cache (location_t loc, c_p
   tree stmt, clauses;
 
   clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
-  clauses = c_finish_omp_clauses (clauses);
+  clauses = c_finish_omp_clauses (clauses, false);
 
   c_parser_skip_to_pragma_eol (parser);
 
@@ -13902,7 +13902,7 @@  omp_split_clauses (location_t loc, enum
   c_omp_split_clauses (loc, code, mask, clauses, cclauses);
   for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
     if (cclauses[i])
-      cclauses[i] = c_finish_omp_clauses (cclauses[i]);
+      cclauses[i] = c_finish_omp_clauses (cclauses[i], true);
 }
 
 /* OpenMP 4.0:
@@ -14668,9 +14668,10 @@  c_parser_omp_target_data (location_t loc
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14800,9 +14801,10 @@  c_parser_omp_target_enter_data (location
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14885,9 +14887,10 @@  c_parser_omp_target_exit_data (location_
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -15016,6 +15019,7 @@  c_parser_omp_target (c_parser *parser, e
 	  TREE_TYPE (stmt) = void_type_node;
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = block;
+	  OMP_TARGET_COMBINED (stmt) = 1;
 	  add_stmt (stmt);
 	  pc = &OMP_TARGET_CLAUSES (stmt);
 	  goto check_clauses;
@@ -15078,7 +15082,7 @@  check_clauses:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -16379,7 +16383,7 @@  c_parser_cilk_for (c_parser *parser, tre
   tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
   OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
   OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
-  clauses = c_finish_omp_clauses (clauses);
+  clauses = c_finish_omp_clauses (clauses, false);
 
   tree block = c_begin_compound_stmt (true);
   tree sb = push_stmt_list ();
@@ -16444,7 +16448,7 @@  c_parser_cilk_for (c_parser *parser, tre
       OMP_CLAUSE_OPERAND (c, 0)
 	= cilk_for_number_of_iterations (omp_for);
       OMP_CLAUSE_CHAIN (c) = clauses;
-      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c);
+      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true);
       add_stmt (omp_par);
     }
 
--- gcc/c/c-typeck.c.jj	2015-07-17 13:06:58.000000000 +0200
+++ gcc/c/c-typeck.c	2015-07-29 16:14:08.276065810 +0200
@@ -11850,7 +11850,7 @@  handle_omp_array_sections_1 (tree c, tre
 /* Handle array sections for clause C.  */
 
 static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
@@ -12030,9 +12030,26 @@  handle_omp_array_sections (tree c)
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 	return false;
       gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
+      if (is_omp)
+	switch (OMP_CLAUSE_MAP_KIND (c))
+	  {
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	    OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+	    break;
+	  default:
+	    break;
+	  }
       tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
-      if (!c_mark_addressable (t))
+      OMP_CLAUSE_SET_MAP_KIND (c2, is_omp
+				   ? GOMP_MAP_FIRSTPRIVATE_POINTER
+				   : GOMP_MAP_POINTER);
+      if (!is_omp && !c_mark_addressable (t))
 	return false;
       OMP_CLAUSE_DECL (c2) = t;
       t = build_fold_addr_expr (first);
@@ -12097,7 +12114,7 @@  c_find_omp_placeholder_r (tree *tp, int
    Remove any elements from the list that are invalid.  */
 
 tree
-c_finish_omp_clauses (tree clauses, bool declare_simd)
+c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
   bitmap_head aligned_head, map_head;
@@ -12136,7 +12153,7 @@  c_finish_omp_clauses (tree clauses, bool
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		{
 		  remove = true;
 		  break;
@@ -12496,7 +12513,7 @@  c_finish_omp_clauses (tree clauses, bool
 	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		remove = true;
 	      break;
 	    }
@@ -12519,7 +12536,7 @@  c_finish_omp_clauses (tree clauses, bool
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		remove = true;
 	      else
 		{
@@ -12556,6 +12573,8 @@  c_finish_omp_clauses (tree clauses, bool
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		     && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 			 || (OMP_CLAUSE_MAP_KIND (c)
+			     == GOMP_MAP_FIRSTPRIVATE_POINTER)
+			 || (OMP_CLAUSE_MAP_KIND (c)
 			     == GOMP_MAP_FORCE_DEVICEPTR)))
 		   && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
 	    {
@@ -12624,10 +12643,11 @@  c_finish_omp_clauses (tree clauses, bool
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  t = OMP_CLAUSE_DECL (c);
-	  if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+	  if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE
+	      && TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
-			"%qs variable is not a pointer",
+			"%qs variable is neither a pointer nor an array",
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
--- gcc/cp/parser.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/cp/parser.c	2015-07-23 12:46:22.172652420 +0200
@@ -32276,27 +32276,28 @@  cp_parser_omp_target_data (cp_parser *pa
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_TO:
-	 case GOMP_MAP_ALWAYS_TO:
-	 case GOMP_MAP_FROM:
-	 case GOMP_MAP_ALWAYS_FROM:
-	 case GOMP_MAP_TOFROM:
-	 case GOMP_MAP_ALWAYS_TOFROM:
-	 case GOMP_MAP_ALLOC:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target data%> with map-type other "
-		     "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
-		     "on %<map%> clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_ALLOC:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target data%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32370,22 +32371,23 @@  cp_parser_omp_target_enter_data (cp_pars
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_TO:
-	 case GOMP_MAP_ALWAYS_TO:
-	 case GOMP_MAP_ALLOC:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target enter data%> with map-type other "
-		     "than %<to%> or %<alloc%> on %<map%> clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_ALLOC:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target enter data%> with map-type other "
+		      "than %<to%> or %<alloc%> on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32455,24 +32457,25 @@  cp_parser_omp_target_exit_data (cp_parse
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_FROM:
-	 case GOMP_MAP_ALWAYS_FROM:
-	 case GOMP_MAP_RELEASE:
-	 case GOMP_MAP_DELETE:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target exit data%> with map-type other "
-		     "than %<from%>, %<release%> or %<delete%> on %<map%>"
-		     " clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_RELEASE:
+	  case GOMP_MAP_DELETE:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target exit data%> with map-type other "
+		      "than %<from%>, %<release%> or %<delete%> on %<map%>"
+		      " clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32637,6 +32640,7 @@  cp_parser_omp_target (cp_parser *parser,
 	  TREE_TYPE (stmt) = void_type_node;
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = body;
+	  OMP_TARGET_COMBINED (stmt) = 1;
 	  add_stmt (stmt);
 	  pc = &OMP_TARGET_CLAUSES (stmt);
 	  goto check_clauses;
@@ -32697,7 +32701,7 @@  check_clauses:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
--- gcc/cp/semantics.c.jj	2015-07-17 13:59:27.000000000 +0200
+++ gcc/cp/semantics.c	2015-07-29 16:14:49.040467753 +0200
@@ -4650,7 +4650,7 @@  handle_omp_array_sections_1 (tree c, tre
 /* Handle array sections for clause C.  */
 
 static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
@@ -4826,10 +4826,26 @@  handle_omp_array_sections (tree c)
 	  OMP_CLAUSE_SIZE (c) = size;
 	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 	    return false;
+	  if (is_omp)
+	    switch (OMP_CLAUSE_MAP_KIND (c))
+	      {
+	      case GOMP_MAP_ALLOC:
+	      case GOMP_MAP_TO:
+	      case GOMP_MAP_FROM:
+	      case GOMP_MAP_TOFROM:
+	      case GOMP_MAP_ALWAYS_TO:
+	      case GOMP_MAP_ALWAYS_FROM:
+	      case GOMP_MAP_ALWAYS_TOFROM:
+		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+		break;
+	      default:
+		break;
+	      }
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
-	  if (!cxx_mark_addressable (t))
+	  OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER
+					      : GOMP_MAP_POINTER);
+	  if (!is_omp && !cxx_mark_addressable (t))
 	    return false;
 	  OMP_CLAUSE_DECL (c2) = t;
 	  t = build_fold_addr_expr (first);
@@ -4847,7 +4863,8 @@  handle_omp_array_sections (tree c)
 	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
 	  OMP_CLAUSE_CHAIN (c) = c2;
 	  ptr = OMP_CLAUSE_DECL (c2);
-	  if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
+	  if (!is_omp
+	      && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
 	      && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
 	    {
 	      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
@@ -5569,7 +5586,7 @@  finish_omp_clauses (tree clauses, bool a
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		{
 		  remove = true;
 		  break;
@@ -6155,7 +6172,7 @@  finish_omp_clauses (tree clauses, bool a
 	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		remove = true;
 	      break;
 	    }
@@ -6189,7 +6206,7 @@  finish_omp_clauses (tree clauses, bool a
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		remove = true;
 	      else
 		{
@@ -6242,7 +6259,9 @@  finish_omp_clauses (tree clauses, bool a
 		   && !cxx_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		     && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
+		     && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+			 || (OMP_CLAUSE_MAP_KIND (c)
+			     == GOMP_MAP_FIRSTPRIVATE_POINTER)))
 		   && !type_dependent_expression_p (t)
 		   && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
 					      == REFERENCE_TYPE)
@@ -6428,12 +6447,14 @@  finish_omp_clauses (tree clauses, bool a
 	    {
 	      tree type = TREE_TYPE (t);
 	      if (TREE_CODE (type) != POINTER_TYPE
+		  && TREE_CODE (type) != ARRAY_TYPE
 		  && (TREE_CODE (type) != REFERENCE_TYPE
-		      || TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE))
+		      || (TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE
+			  && TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
-			    "%qs variable is not a pointer or reference "
-			    "to pointer",
+			    "%qs variable is neither a pointer, nor an array"
+			    "nor reference to pointer or array",
 			    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 		  remove = true;
 		}
--- include/gomp-constants.h.jj	2015-07-21 09:07:23.689851239 +0200
+++ include/gomp-constants.h	2015-07-29 16:15:20.101012063 +0200
@@ -74,6 +74,17 @@  enum gomp_map_kind
     GOMP_MAP_FORCE_DEVICEPTR =		(GOMP_MAP_FLAG_SPECIAL_1 | 0),
     /* Do not map, copy bits for firstprivate instead.  */
     GOMP_MAP_FIRSTPRIVATE =		(GOMP_MAP_FLAG_SPECIAL | 0),
+    /* Similarly, but store the value in the pointer rather than
+       pointed by the pointer.  */
+    GOMP_MAP_FIRSTPRIVATE_INT =		(GOMP_MAP_FLAG_SPECIAL | 1),
+    /* Pointer translate host address into device address and copy that
+       back to host.  */
+    GOMP_MAP_USE_DEVICE_PTR =		(GOMP_MAP_FLAG_SPECIAL | 2),
+    /* Allocate a zero length array section.  Prefer next non-zero length
+       mapping over previous non-zero length mapping over zero length mapping
+       at the address.  If not already mapped, do nothing (and pointer translate
+       to NULL).  */
+    GOMP_MAP_ZERO_LEN_ARRAY_SECTION = 	(GOMP_MAP_FLAG_SPECIAL | 3),
     /* Allocate.  */
     GOMP_MAP_FORCE_ALLOC =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
     /* ..., and copy to device.  */
@@ -95,7 +106,11 @@  enum gomp_map_kind
     GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
-					 | GOMP_MAP_FORCE_DEALLOC)
+					 | GOMP_MAP_FORCE_DEALLOC),
+
+    /* Internal to GCC, not used in libgomp.  */
+    /* Do not map, but pointer assign a pointer instead.  */
+    GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
--- libgomp/libgomp.h.jj	2015-07-15 13:00:32.000000000 +0200
+++ libgomp/libgomp.h	2015-07-22 21:09:39.023307107 +0200
@@ -647,11 +647,9 @@  struct target_var_desc {
   bool copy_from;
   /* True if data always should be copied from device to host at the end.  */
   bool always_copy_from;
-  /* Used for unmapping of array sections, can be nonzero only when
-     always_copy_from is true.  */
+  /* Relative offset against key host_start.  */
   uintptr_t offset;
-  /* Used for unmapping of array sections, can be less than the size of the
-     whole object only when always_copy_from is true.  */
+  /* Actual length.  */
   uintptr_t length;
 };
 
--- libgomp/target.c.jj	2015-07-21 09:07:23.690851224 +0200
+++ libgomp/target.c	2015-07-29 17:12:06.377060519 +0200
@@ -142,7 +142,26 @@  resolve_device (int device_id)
 }
 
 
-/* Handle the case where splay_tree_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_lookup (splay_tree mem_map, splay_tree_key key)
+{
+  if (key->host_start != key->host_end)
+    return splay_tree_lookup (mem_map, key);
+
+  key->host_end++;
+  splay_tree_key n = splay_tree_lookup (mem_map, key);
+  key->host_end--;
+  if (n)
+    return n;
+  key->host_start--;
+  n = splay_tree_lookup (mem_map, key);
+  key->host_start++;
+  if (n)
+    return n;
+  return splay_tree_lookup (mem_map, key);
+}
+
+/* Handle the case where gomp_map_lookup found oldn for newn.
    Helper function of gomp_map_vars.  */
 
 static inline void
@@ -204,20 +223,8 @@  gomp_map_pointer (struct target_mem_desc
     }
   /* Add bias to the pointer value.  */
   cur_node.host_start += bias;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      /* Could be possibly zero size array section.  */
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-      if (n == NULL)
-	{
-	  cur_node.host_start--;
-	  n = splay_tree_lookup (mem_map, &cur_node);
-	  cur_node.host_start++;
-	}
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n == NULL)
     {
       gomp_mutex_unlock (&devicep->lock);
@@ -271,9 +278,29 @@  gomp_map_vars (struct gomp_device_descr
   for (i = 0; i < mapnum; i++)
     {
       int kind = get_kind (short_mapkind, kinds, i);
-      if (hostaddrs[i] == NULL)
+      if (hostaddrs[i] == NULL
+	  || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
 	{
 	  tgt->list[i].key = NULL;
+	  tgt->list[i].offset = ~(uintptr_t) 0;
+	  continue;
+	}
+      else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
+	{
+	  cur_node.host_start = (uintptr_t) hostaddrs[i];
+	  cur_node.host_end = cur_node.host_start;
+	  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
+	  if (n == NULL)
+	    {
+	      gomp_mutex_unlock (&devicep->lock);
+	      gomp_fatal ("use_device_ptr pointer wasn't mapped");
+	    }
+	  cur_node.host_start -= n->host_start;
+	  hostaddrs[i]
+	    = (void *) (n->tgt->tgt_start + n->tgt_offset
+			+ cur_node.host_start);
+	  tgt->list[i].key = NULL;
+	  tgt->list[i].offset = ~(uintptr_t) 0;
 	  continue;
 	}
       cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -293,7 +320,19 @@  gomp_map_vars (struct gomp_device_descr
 	  has_firstprivate = true;
 	  continue;
 	}
-      splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+      splay_tree_key n;
+      if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
+	{
+	  n = gomp_map_lookup (mem_map, &cur_node);
+	  if (!n)
+	    {
+	      tgt->list[i].key = NULL;
+	      tgt->list[i].offset = ~(uintptr_t) 1;
+	      continue;
+	    }
+	}
+      else
+	n = splay_tree_lookup (mem_map, &cur_node);
       if (n)
 	gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
 				kind & typemask);
@@ -386,6 +425,15 @@  gomp_map_vars (struct gomp_device_descr
 		tgt_size += len;
 		continue;
 	      }
+	    switch (kind & typemask)
+	      {
+	      case GOMP_MAP_FIRSTPRIVATE_INT:
+	      case GOMP_MAP_USE_DEVICE_PTR:
+	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+		continue;
+	      default:
+		break;
+	      }
 	    splay_tree_key k = &array->key;
 	    k->host_start = (uintptr_t) hostaddrs[i];
 	    if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -518,15 +566,18 @@  gomp_map_vars (struct gomp_device_descr
 	{
 	  if (tgt->list[i].key == NULL)
 	    {
-	      if (hostaddrs[i] == NULL)
-		cur_node.tgt_offset = (uintptr_t) NULL;
+	      if (tgt->list[i].offset == ~(uintptr_t) 0)
+		cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
+	      else if (tgt->list[i].offset == ~(uintptr_t) 1)
+		cur_node.tgt_offset = 0;
 	      else
 		cur_node.tgt_offset = tgt->tgt_start
 				      + tgt->list[i].offset;
 	    }
 	  else
 	    cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
-				  + tgt->list[i].key->tgt_offset;
+				  + tgt->list[i].key->tgt_offset
+				  + tgt->list[i].offset;
 	  /* FIXME: see above FIXME comment.  */
 	  devicep->host2dev_func (devicep->target_id,
 				  (void *) (tgt->tgt_start
@@ -1052,7 +1103,38 @@  GOMP_target_41 (int device, void (*fn) (
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
-    return gomp_target_fallback (fn, hostaddrs);
+    {
+      size_t i, tgt_align = 0, tgt_size = 0;
+      char *tgt = NULL;
+      for (i = 0; i < mapnum; i++)
+	if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+	  {
+	    size_t align = (size_t) 1 << (kinds[i] >> 8);
+	    if (tgt_align < align)
+	      tgt_align = align;
+	    tgt_size = (tgt_size + align - 1) & ~(align - 1);
+	    tgt_size += sizes[i];
+	  }
+      if (tgt_align)
+	{
+	  tgt = gomp_alloca (tgt_size + tgt_align - 1);
+	  uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
+	  if (al)
+	    tgt += tgt_align - al;
+	  tgt_size = 0;
+	  for (i = 0; i < mapnum; i++)
+	    if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+	      {
+		size_t align = (size_t) 1 << (kinds[i] >> 8);
+		tgt_size = (tgt_size + align - 1) & ~(align - 1);
+		memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
+		hostaddrs[i] = tgt + tgt_size;
+		tgt_size = tgt_size + sizes[i];
+	      }
+	}
+      gomp_target_fallback (fn, hostaddrs);
+      return;
+    }
 
   void *fn_addr = gomp_get_target_fn_addr (devicep, fn);
 
@@ -1289,20 +1371,8 @@  omp_target_is_present (void *ptr, size_t
   struct splay_tree_key_s cur_node;
 
   cur_node.host_start = (uintptr_t) ptr + offset;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      /* Could be possibly zero size array section.  */
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-      if (n == NULL)
-	{
-	  cur_node.host_start--;
-	  n = splay_tree_lookup (mem_map, &cur_node);
-	  cur_node.host_start++;
-	}
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   int ret = n != NULL;
   gomp_mutex_unlock (&devicep->lock);
   return ret;
@@ -1524,7 +1594,7 @@  omp_target_associate_ptr (void *host_ptr
 
   cur_node.host_start = (uintptr_t) host_ptr;
   cur_node.host_end = cur_node.host_start + size;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n)
     {
       if (n->tgt->tgt_start + n->tgt_offset
@@ -1584,13 +1654,8 @@  omp_target_disassociate_ptr (void *ptr,
   int ret = EINVAL;
 
   cur_node.host_start = (uintptr_t) ptr;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n
       && n->host_start == cur_node.host_start
       && n->refcount == REFCOUNT_INFINITY
--- libgomp/testsuite/libgomp.c++/target-2.C.jj	2015-06-30 14:24:03.000000000 +0200
+++ libgomp/testsuite/libgomp.c++/target-2.C	2015-07-23 17:48:08.978674497 +0200
@@ -33,7 +33,8 @@  fn2 (int x, double (&dr) [1024], double
   int j;
   fn1 (hr + 2 * x, ir + 2 * x, x);
   #pragma omp target map(to: br[:x], cr[0:x], dr[x:x], er[x:x]) \
-		     map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x])
+		     map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x]) \
+		     map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (j = 0; j < x; j++)
 	s += br[j] * cr[j] + dr[x + j] + er[x + j]
--- libgomp/testsuite/libgomp.c++/target-7.C.jj	2015-07-22 11:36:53.042867520 +0200
+++ libgomp/testsuite/libgomp.c++/target-7.C	2015-07-22 11:32:00.000000000 +0200
@@ -0,0 +1,90 @@ 
+extern "C" void abort ();
+
+void
+foo (int *x, int *&y, int (&z)[15])
+{
+  int a[10], b[15], err, i;
+  for (i = 0; i < 10; i++)
+    a[i] = 7 * i;
+  for (i = 0; i < 15; i++)
+    b[i] = 8 * i;
+  #pragma omp target map(to:x[5:10], y[5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if (x[5 + i] != 20 + 4 * i
+	  || y[5 + i] != 25 + 5 * i
+	  || z[5 + i] != 30 + 6 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n, int v)
+{
+  int a[n], b[n], c[n], d[n], e[n], err, i;
+  int (*x)[n] = &c;
+  int (*y2)[n] = &d;
+  int (*&y)[n] = y2;
+  int (&z)[n] = e;
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 4 * i;
+      (*y)[i] = 5 * i;
+      z[i] = 6 * i;
+      a[i] = 7 * i;
+      b[i] = 8 * i;
+    }
+  #pragma omp target map(to:x[0][5:10], y[0][5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 20 + 4 * i
+	  || (*y)[5 + i] != 25 + 5 * i
+	  || z[5 + i] != 30 + 6 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 9 * i;
+      (*y)[i] = 10 * i;
+      z[i] = 11 * i;
+      a[i] = 12 * i;
+      b[i] = 13 * i;
+    }
+  #pragma omp target map(to:x[0][v:v+5], y[0][v:v+5], z[v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 45 + 9 * i
+	  || (*y)[5 + i] != 50 + 10 * i
+	  || z[5 + i] != 55 + 11 * i
+	  || a[i] != 12 * i
+	  || b[5 + i] != 65 + 13 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int x[15], y2[15], z[15], *y = y2, i;
+  for (i = 0; i < 15; i++)
+    {
+      x[i] = 4 * i;
+      y[i] = 5 * i;
+      z[i] = 6 * i;
+    }
+  foo (x, y, z);
+  bar (15, 5);
+}
--- libgomp/testsuite/libgomp.c++/target-8.C.jj	2015-07-27 13:39:49.446401028 +0200
+++ libgomp/testsuite/libgomp.c++/target-8.C	2015-07-27 13:39:27.000000000 +0200
@@ -0,0 +1,58 @@ 
+extern "C" void abort ();
+struct S { int a; };
+#ifdef __SIZEOF_INT128__
+typedef __int128 T;
+#else
+typedef long long int T;
+#endif
+
+void
+foo (T a, int b, struct S c)
+{
+  int err;
+  #pragma omp target firstprivate (a, b, c) map(from:err)
+  {
+    err = 0;
+    if (a != 131 || b != 276 || c.a != 59)
+      err = 1;
+    a = 936;
+    b = 27;
+    c.a = 98;
+    if (a != 936 || b != 27 || c.a != 98)
+      err = 1;
+  }
+  if (err || a != 131 || b != 276 || c.a != 59)
+    abort ();
+}
+
+void
+bar (T &a, int &b, struct S &c)
+{
+  int err;
+  #pragma omp target firstprivate (a, b, c) map(from:err)
+  {
+    err = 0;
+    if (a != 131 || b != 276 || c.a != 59)
+      err = 1;
+    a = 936;
+    b = 27;
+    c.a = 98;
+    if (a != 936 || b != 27 || c.a != 98)
+      err = 1;
+  }
+  if (err || a != 131 || b != 276 || c.a != 59)
+    abort ();
+}
+
+int
+main ()
+{
+  T a = 131;
+  int b = 276;
+  struct S c;
+  c.a = 59;
+  foo (a, b, c);
+  bar (a, b, c);
+  if (a != 131 || b != 276 || c.a != 59)
+    abort ();
+}
--- libgomp/testsuite/libgomp.c++/target-9.C.jj	2015-07-28 16:57:29.940191999 +0200
+++ libgomp/testsuite/libgomp.c++/target-9.C	2015-07-28 20:30:05.951617430 +0200
@@ -0,0 +1,73 @@ 
+extern "C" void abort (void);
+
+void
+foo (int *&p, int (&s)[5], int n)
+{
+  int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
+  int *r = a + 1, *q = p - 1, i, err;
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(r) map(from:err)
+  #pragma omp target is_device_ptr(r) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (r[i - 1] != 7 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:q[:4])
+  #pragma omp target data use_device_ptr(p) map(from:err)
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:b)
+  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:c)
+  #pragma omp target data use_device_ptr(c) map(from:err)
+  #pragma omp target is_device_ptr(c) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 3; i++)
+      if (c[i] != 20 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:s[:5])
+  #pragma omp target data use_device_ptr(s) map(from:err)
+  #pragma omp target is_device_ptr(s) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 5; i++)
+      if (s[i] != 17 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
+  int *p = a + 1;
+  foo (p, b, 9);
+}
--- libgomp/testsuite/libgomp.c/target-1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-1.c	2015-07-23 17:08:32.474133124 +0200
@@ -34,7 +34,7 @@  fn2 (int x, int y, int z)
   fn1 (b, c, x);
   #pragma omp target data map(to: b)
   {
-    #pragma omp target map(tofrom: c)
+    #pragma omp target map(tofrom: c, s)
       #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)
@@ -52,7 +52,7 @@  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 target map(to: b, c) map(tofrom:s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	tgt (), s += b[i] * c[i];
@@ -66,7 +66,8 @@  fn4 (int x, double *p)
   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 target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+		     map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + p[x + i];
--- libgomp/testsuite/libgomp.c/target-2.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-2.c	2015-07-23 17:09:27.987350372 +0200
@@ -23,7 +23,7 @@  fn2 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d + x, x);
-  #pragma omp target map(to: b, c[:x], d[x:x], e)
+  #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
@@ -38,7 +38,7 @@  fn3 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d, x);
-  #pragma omp target
+  #pragma omp target map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[i];
@@ -56,7 +56,7 @@  fn4 (int x)
   #pragma omp target data map(from: b, c[:x], d[x:x], e)
     {
       #pragma omp target update to(b, c[:x], d[x:x], e)
-      #pragma omp target map(c[:x], d[x:x])
+      #pragma omp target map(c[:x], d[x:x], s)
 	#pragma omp parallel for reduction(+:s)
 	  for (i = 0; i < x; i++)
 	    {
--- libgomp/testsuite/libgomp.c/target-7.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-7.c	2015-07-23 17:12:33.159753962 +0200
@@ -37,63 +37,63 @@  foo (int f)
     abort ();
   #pragma omp target data device (d) map (to: h)
   {
-    #pragma omp target device (d)
+    #pragma omp target device (d) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
       abort ();
     #pragma omp target update device (d) from (h)
   }
   #pragma omp target data if (v > 1) map (to: h)
   {
-    #pragma omp target if (v > 1)
+    #pragma omp target if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
       abort ();
     #pragma omp target update if (v > 1) from (h)
   }
   #pragma omp target data device (d) if (v > 1) map (to: h)
   {
-    #pragma omp target device (d) if (v > 1)
+    #pragma omp target device (d) if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
       abort ();
     #pragma omp target update device (d) if (v > 1) from (h)
   }
   #pragma omp target data if (v <= 1) map (to: h)
   {
-    #pragma omp target if (v <= 1)
+    #pragma omp target if (v <= 1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 8)
       abort ();
     #pragma omp target update if (v <= 1) from (h)
   }
   #pragma omp target data device (d) if (v <= 1) map (to: h)
   {
-    #pragma omp target device (d) if (v <= 1)
+    #pragma omp target device (d) if (v <= 1) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
       abort ();
     #pragma omp target update device (d) if (v <= 1) from (h)
   }
   #pragma omp target data if (0) map (to: h)
   {
-    #pragma omp target if (0)
+    #pragma omp target if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
       abort ();
     #pragma omp target update if (0) from (h)
   }
   #pragma omp target data device (d) if (0) map (to: h)
   {
-    #pragma omp target device (d) if (0)
+    #pragma omp target device (d) if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
       abort ();
     #pragma omp target update device (d) if (0) from (h)
   }
   #pragma omp target data if (1) map (to: h)
   {
-    #pragma omp target if (1)
+    #pragma omp target if (1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 12)
       abort ();
     #pragma omp target update if (1) from (h)
   }
   #pragma omp target data device (d) if (1) map (to: h)
   {
-    #pragma omp target device (d) if (1)
+    #pragma omp target device (d) if (1) map (tofrom: h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
       abort ();
     #pragma omp target update device (d) if (1) from (h)
--- libgomp/testsuite/libgomp.c/target-15.c.jj	2015-07-22 11:37:11.655612690 +0200
+++ libgomp/testsuite/libgomp.c/target-15.c	2015-07-23 21:53:37.354632916 +0200
@@ -0,0 +1,74 @@ 
+extern void abort (void);
+
+void
+foo (int *x)
+{
+  int a[10], b[15], err, i;
+  for (i = 0; i < 10; i++)
+    a[i] = 7 * i;
+  for (i = 0; i < 15; i++)
+    b[i] = 8 * i;
+  #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if (x[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n, int v)
+{
+  int a[n], b[n], c[n], d[n], e[n], err, i;
+  int (*x)[n] = &c;
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 4 * i;
+      a[i] = 7 * i;
+      b[i] = 8 * i;
+    }
+  #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 9 * i;
+      a[i] = 12 * i;
+      b[i] = 13 * i;
+    }
+  #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 45 + 9 * i
+	  || a[i] != 12 * i
+	  || b[5 + i] != 65 + 13 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int x[15], i;
+  for (i = 0; i < 15; i++)
+    x[i] = 4 * i;
+  foo (x);
+  bar (15, 5);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/target-16.c.jj	2015-07-23 21:53:28.905753778 +0200
+++ libgomp/testsuite/libgomp.c/target-16.c	2015-07-24 12:20:32.048722516 +0200
@@ -0,0 +1,45 @@ 
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 7 * i;
+  #pragma omp target firstprivate (a) map(from:err) private (i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n)
+{
+  int a[n], i, err;
+  #pragma omp target private (a) map(from:err)
+  {
+    #pragma omp parallel for
+    for (i = 0; i < n; i++)
+      a[i] = 7 * i;
+    err = 0;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err |= 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (7);
+  bar (7);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/target-17.c.jj	2015-07-24 19:50:14.275109272 +0200
+++ libgomp/testsuite/libgomp.c/target-17.c	2015-07-24 19:47:57.000000000 +0200
@@ -0,0 +1,99 @@ 
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 5 * i;
+  #pragma omp target map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 5 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 6 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  int on = n;
+  #pragma omp target firstprivate (n) map(tofrom: n)
+  {
+    n++;
+  }
+  if (on != n)
+    abort ();
+  #pragma omp target map(tofrom: n) private (n)
+  {
+    n = 25;
+  }
+  if (on != n)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 9 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(tofrom:a) map(from:err) private(a, i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      a[i] = 7;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7)
+	err |= 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    if (a[i] != 10 * i)
+      abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/target-18.c.jj	2015-07-28 16:50:12.139587099 +0200
+++ libgomp/testsuite/libgomp.c/target-18.c	2015-07-28 19:59:41.000000000 +0200
@@ -0,0 +1,52 @@ 
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[4] = { 0, 1, 2, 3 }, b[n];
+  int *p = a + 1, i, err;
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(p) map(from:err)
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < 4; i++)
+    a[i] = 23 + i;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(a) map(from:err)
+  #pragma omp target is_device_ptr(a) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (a[i] != 23 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:b)
+  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/target-19.c.jj	2015-07-29 16:28:01.783837512 +0200
+++ libgomp/testsuite/libgomp.c/target-19.c	2015-07-29 16:32:42.800714833 +0200
@@ -0,0 +1,127 @@ 
+extern void abort (void);
+
+void
+foo (int *p, int *q, int *r, int n, int m)
+{
+  int i, err, *s = r;
+  #pragma omp target data map(to:p[0:8])
+  {
+    /* For zero length array sections, p points to the start of
+       already mapped range, q to the end of it, and r does not point
+       to an mapped range.  */
+    #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* Implicit mapping of pointers behaves the same way.  */
+    #pragma omp target private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* And zero-length array sections, though not known at compile
+       time, behave the same.  */
+    #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1 || q[i - 8] != i + 1)
+	  err = 1;
+      if (p + 8 != q || (r != (int *) 0 && r != s))
+	err = 1;
+    }
+    if (err)
+      abort ();
+    /* Non-zero length array sections, though not known at compile,
+       behave differently.  */
+    #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err)
+    {
+      err = 0;
+      for (i = 0; i < 8; i++)
+	if (p[i] != i + 1)
+	  err = 1;
+      if (q[0] != 9 || r[0] != 10)
+	err = 1;
+    }
+    if (err)
+      abort ();
+    #pragma omp target data map(to:q[0:1])
+    {
+      /* For zero length array sections, p points to the start of
+	 already mapped range, q points to the start of another one,
+	 and r to the end of the second one.  */
+      #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* Implicit mapping of pointers behaves the same way.  */
+      #pragma omp target private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* And zero-length array sections, though not known at compile
+	 time, behave the same.  */
+      #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r != q + 1)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+      /* Non-zero length array sections, though not known at compile,
+	 behave differently.  */
+      #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err)
+      {
+	err = 0;
+	for (i = 0; i < 8; i++)
+	  if (p[i] != i + 1)
+	    err = 1;
+	if (q[0] != 9 || r[0] != 10)
+	  err = 1;
+      }
+      if (err)
+	abort ();
+    }
+  }
+}
+
+int
+main ()
+{
+  int a[32], i;
+  for (i = 0; i < 32; i++)
+    a[i] = i;
+  foo (a + 1, a + 9, a + 10, 0, 1);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/examples-4/e.51.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.51.3.c	2015-07-23 15:58:15.867779262 +0200
@@ -47,7 +47,7 @@  void gramSchmidt (int Q[][COLS], const i
       {
 	int tmp = 0;
 
-	#pragma omp target
+	#pragma omp target map(tofrom:tmp)
 	  #pragma omp parallel for reduction(+:tmp)
 	    for (i = 0; i < rows; i++)
 	      tmp += (Q[i][k] * Q[i][k]);
--- libgomp/testsuite/libgomp.c/examples-4/e.53.1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.1.c	2015-07-23 15:59:44.430518114 +0200
@@ -20,7 +20,7 @@  int fib_wrapper (int n)
 {
   int x = 0;
 
-  #pragma omp target if(n > THRESHOLD)
+  #pragma omp target if(n > THRESHOLD) map(from:x)
     x = fib (n);
 
   return x;
--- libgomp/testsuite/libgomp.c/examples-4/e.53.4.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.4.c	2015-07-23 16:00:22.468976440 +0200
@@ -41,7 +41,7 @@  float accum (int k)
   int i;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	tmp += Pfun (i, k);
--- libgomp/testsuite/libgomp.c/examples-4/e.53.5.c.jj	2015-06-17 21:00:36.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.5.c	2015-07-23 16:01:17.802188485 +0200
@@ -48,7 +48,7 @@  float accum ()
   int i, k;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	{
--- libgomp/testsuite/libgomp.c/examples-4/e.54.2.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.2.c	2015-07-23 16:02:02.343554209 +0200
@@ -32,7 +32,7 @@  float dotprod (float B[], float C[], int
   int i, i0;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp teams num_teams(num_teams) thread_limit(block_threads) \
 		      reduction(+:sum)
       #pragma omp distribute
--- libgomp/testsuite/libgomp.c/examples-4/e.54.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.3.c	2015-07-23 16:02:28.060187999 +0200
@@ -31,7 +31,7 @@  float dotprod (float B[], float C[], int
   int i;
   float sum = 0;
 
-  #pragma omp target teams map(to: B[0:n], C[0:n])
+  #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp distribute parallel for reduction(+:sum)
       for (i = 0; i < n; i++)
 	sum += B[i] * C[i];
--- libgomp/testsuite/libgomp.c/examples-4/e.54.4.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.4.c	2015-07-23 16:03:21.446427770 +0200
@@ -31,7 +31,7 @@  float dotprod (float B[], float C[], int
   int i;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum)
     #pragma omp teams num_teams(8) thread_limit(16)
       #pragma omp distribute parallel for reduction(+:sum) \
 					  dist_schedule(static, 1024) \
--- libgomp/testsuite/libgomp.c/examples-4/e.57.1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.57.1.c	2015-07-23 17:37:01.880139916 +0200
@@ -10,11 +10,11 @@  int main ()
   int b = 0;
   int c, d;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -26,11 +26,11 @@  int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -42,11 +42,11 @@  int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
--- libgomp/testsuite/libgomp.c/examples-4/e.57.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.57.3.c	2015-07-23 16:08:48.176775074 +0200
@@ -9,7 +9,7 @@  int main ()
   int res;
   int default_device = omp_get_default_device ();
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (res)
@@ -17,7 +17,7 @@  int main ()
 
   omp_set_default_device (omp_get_num_devices ());
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (!res)