[OpenMP,5.0] Improve OpenMP target support for C++ [PR92120 v4]

Message ID c9454ea4-c5d6-f210-f08b-914ed7b46859@codesourcery.com
State New
Headers show
Series
  • [OpenMP,5.0] Improve OpenMP target support for C++ [PR92120 v4]
Related show

Commit Message

Chung-Lin Tang June 18, 2021, 2:25 p.m.
Hi Jakub,
this patch is the "v4" version of my PR92120 patch, v3 was here:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570886.html

(there I listed the various patches from devel/omp/gcc-10 branch that was combined,
which I won't repeat here).

Basically this v4 adds fixes for lambda capture, which was already pushed to
  devel/omp/gcc-11 yesterday:
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/572988.html

I have attached both the combined v4 version, and the v3-to-v4 diff.

Tested on x86_64-linux with nvptx offloading, seeking for approval to trunk.

Thanks,
Chung-Lin

	gcc/cp/
	* cp-tree.h (finish_omp_target): New declaration.
	(finish_omp_target_clauses): Likewise.
	* parser.c (cp_parser_omp_clause_map): Adjust call to
	cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true.
	(cp_parser_omp_target): Factor out code, adjust into calls to new
	function finish_omp_target.
	* pt.c (tsubst_expr): Add call to finish_omp_target_clauses for
	OMP_TARGET case.
	* semantics.c (handle_omp_array_sections_1): Add handling to create
	'this->member' from 'member' FIELD_DECL.
	(handle_omp_array_sections): Likewise.
	(finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP
	map clauses. Handle 'A->member' case in map clauses.
	(struct omp_target_walk_data): New struct for walking over
	target-directive tree body.
	(finish_omp_target_clauses_r): New function for tree walk.
	(finish_omp_target_clauses): New function.
	(finish_omp_target): New function.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in
	call to c_parser_omp_variable_list to 'true'.
	* c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in
	array base handling.
	(c_finish_omp_clauses): Handle 'A->member' case in map clauses.

	gcc/
	* gimplify.c ("tree-hash-traits.h"): Add include.
	(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
	hash_map<tree_operand, tree> *. Adjust struct map handling to handle
	cases of *A and A->B expressions. Under !DECL_P case of
	GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to
	struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when
	handling component_ref_p case. Add unshare_expr and gimplification
	when created GOMP_MAP_STRUCT is not a DECL. Add code to add
	firstprivate pointer for *pointer-to-struct case.
	(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
	exit data directives code to earlier position.
	* omp-low.c (lower_omp_target):
	Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
	GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
	* tree-pretty-print.c (dump_omp_clause): Likewise.

	gcc/testsuite/
	* gcc.dg/gomp/target-3.c: New testcase.
	* g++.dg/gomp/target-3.C: New testcase.
	* g++.dg/gomp/target-lambda-1.C: New testcase.
	* g++.dg/gomp/target-lambda-2.C: New testcase.
	* g++.dg/gomp/target-this-1.C: New testcase.
	* g++.dg/gomp/target-this-2.C: New testcase.
	* g++.dg/gomp/target-this-3.C: New testcase.
	* g++.dg/gomp/target-this-4.C: New testcase.
	* g++.dg/gomp/target-this-5.C: New testcase.
	* g++.dg/gomp/this-2.C: Adjust testcase.

	include/
	* gomp-constants.h (enum gomp_map_kind):
	Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
	GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
	(GOMP_MAP_POINTER_P):
	Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION.

	libgomp/
	* libgomp.h (gomp_attach_pointer): Add bool parameter.
	* oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer.
	(goacc_enter_data_internal): Likewise.
	* target.c (gomp_map_vars_existing): Update assert condition to
	include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION.
	(gomp_map_pointer): Add 'bool allow_zero_length_array_sections'
	parameter, add support for mapping a pointer with NULL target.
	(gomp_attach_pointer): Add 'bool allow_zero_length_array_sections'
	parameter, add support for attaching a pointer with NULL target.
	(gomp_map_vars_internal): Update calls to gomp_map_pointer and
	gomp_attach_pointer, add handling for
	GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
	GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases.
	* testsuite/libgomp.c++/target-23.C: New testcase.
	* testsuite/libgomp.c++/target-lambda-1.C: New testcase.
	* testsuite/libgomp.c++/target-lambda-2.C: New testcase.
	* testsuite/libgomp.c++/target-this-1.C: New testcase.
	* testsuite/libgomp.c++/target-this-2.C: New testcase.
	* testsuite/libgomp.c++/target-this-3.C: New testcase.
	* testsuite/libgomp.c++/target-this-4.C: New testcase.
	* testsuite/libgomp.c++/target-this-5.C: New testcase.

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 41ddea405ea..455620c0663 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -9049,6 +9049,8 @@ struct omp_target_walk_data
 
   tree current_closure;
   hash_set<tree> closure_vars_accessed;
+
+  hash_set<tree> local_decls;
 };
 
 static tree
@@ -9107,12 +9109,25 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
       return NULL_TREE;
     }
 
+  if (TREE_CODE (t) == BIND_EXPR)
+    {
+      tree block = BIND_EXPR_BLOCK (t);
+      for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var))
+	if (!data->local_decls.contains (var))
+	  data->local_decls.add (var);
+      return NULL_TREE;
+    }
+
   if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t)))
     {
       tree lt = TREE_TYPE (t);
       gcc_assert (CLASS_TYPE_P (lt));
 
-      if (!data->lambda_objects_accessed.contains (t))
+      if (!data->lambda_objects_accessed.contains (t)
+	  /* Do not prepare to create target maps for locally declared
+	     lambdas or anonymous ones.  */
+	  && !data->local_decls.contains (t)
+	  && TREE_CODE (t) != TARGET_EXPR)
 	data->lambda_objects_accessed.add (t);
       *walk_subtrees = 0;
       return NULL_TREE;
@@ -9398,6 +9413,9 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
 	   i != data.lambda_objects_accessed.end (); ++i)
 	{
 	  tree lobj = *i;
+	  if (TREE_CODE (lobj) == TARGET_EXPR)
+	    lobj = TREE_OPERAND (lobj, 0);
+
 	  tree lt = TREE_TYPE (lobj);
 	  gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt));
 
@@ -9434,7 +9452,7 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
 		  tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
 				     lobj, fld, NULL_TREE);
 		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
-		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
 		  OMP_CLAUSE_DECL (c)
 		    = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp);
 		  OMP_CLAUSE_SIZE (c)
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
new file mode 100644
index 00000000000..bdf2564cd04
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
@@ -0,0 +1,35 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}
+
+/* Make sure lambda objects do NOT appear in target maps.  */
+/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
new file mode 100644
index 00000000000..1d3561ffbd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
@@ -0,0 +1,30 @@
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}

Comments

Richard Biener via Gcc-patches June 24, 2021, 1:15 p.m. | #1
On Fri, Jun 18, 2021 at 10:25:16PM +0800, Chung-Lin Tang wrote:

Note, you'll need to rebase your patch, it clashes with
r12-1768-g7619d33471c10fe3d149dcbb701d99ed3dd23528.
Sorry for that.  And sorry for patch review delay.

> --- a/gcc/c/c-typeck.c

> +++ b/gcc/c/c-typeck.c

> @@ -13104,6 +13104,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,

>  		  return error_mark_node;

>  		}

>  	      t = TREE_OPERAND (t, 0);

> +	      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)


Map clauses never appear on declare simd, so
(ort == C_ORT_ACC || ort == C_ORT_OMP)
previously meant always and since the in_reduction change is incorrect
(as C_ORT_OMP_TARGET is used for target construct but not for
e.g. target data* or target update).

> +		  && TREE_CODE (t) == MEM_REF)


So please just use if (TREE_CODE (t) == MEM_REF)
or explain when it shouldn't trigger.

> @@ -14736,6 +14743,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)

>  		    {

>  		      while (TREE_CODE (t) == COMPONENT_REF)

>  			t = TREE_OPERAND (t, 0);

> +		      if (TREE_CODE (t) == MEM_REF)

> +			{

> +			  t = TREE_OPERAND (t, 0);

> +			  STRIP_NOPS (t);

> +			}


This doesn't look correct.  At least the parsing (and the spec AFAIK)
doesn't ensure that if there is ->, it must come before all the dots.
So, if one uses map (s->x.y) the above would work, but if map (s->x.y->z) or
map (s.a->b->c->d->e) is used, it wouldn't.  I'd expect a single
while loop that looks through COMPONENT_REFs and MEM_REFs as they appear.
Maybe the handle_omp_array_sections_1 MEM_REF case too?

Or do you want to have it done incrementally, start with supporting only
a single -> first before all the dots and later on add support for the rest?

I think the 5.0 and especially 5.1 wording basically says that map clause
operand is arbitrary lvalue expression that includes array section support
too, so eventually we should just have somewhere in parsing scope a bool
whether OpenMP array sections are allowed or not, add OMP_ARRAY_REF or
similar tree code for those and after parsing the expression, ensure
array sections appear only where they can appear and for a subset of the
lvalue expressions where we have decl plus series of -> field or . field
or [ index ] or [ array section stuff ] handle those specially.
That arbitrary lvalue can certainly be done incrementally.
map (foo(123)->a.b[3]->c.d[:7]) and the like.

>  		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP

>  			  && OMP_CLAUSE_MAP_IMPLICIT (c)

>  			  && (bitmap_bit_p (&map_head, DECL_UID (t))

> @@ -14802,6 +14814,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)

>  	       bias) to zero here, so it is not set erroneously to the pointer

>  	       size later on in gimplify.c.  */

>  	    OMP_CLAUSE_SIZE (c) = size_zero_node;

> +	  indir_component_ref_p = false;

> +	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)


Same comment about ort tests.

> +	      && TREE_CODE (t) == COMPONENT_REF

> +	      && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)

> +	    {

> +	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);

> +	      indir_component_ref_p = true;

> +	      STRIP_NOPS (t);

> +	    }


Again, this can handle only a single ->

> @@ -42330,16 +42328,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,

>  		    cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;

>  		  }

>  	    }

> -	  tree stmt = make_node (OMP_TARGET);

> -	  TREE_TYPE (stmt) = void_type_node;

> -	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];

> -	  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);

> -	  OMP_TARGET_BODY (stmt) = body;

> -	  OMP_TARGET_COMBINED (stmt) = 1;

> -	  SET_EXPR_LOCATION (stmt, pragma_tok->location);

> -	  add_stmt (stmt);

> -	  pc = &OMP_TARGET_CLAUSES (stmt);

> -	  goto check_clauses;

> +	  c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true);

> +	  finish_omp_target (pragma_tok->location,

> +			     cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);


What is the advantage of finish_omp_target.  Perhaps the check_clauses label
can be renamed and more things common to both paths moved after the label if
needed, but as long as it isn't something also called during instantiation,
I find it cleaner to do it in cp_parser_omp_target at one place.
The reason for e.g. finish_omp_parallel is that it is called from both
parsing and instantiation.

> --- a/gcc/cp/semantics.c

> +++ b/gcc/cp/semantics.c

> @@ -4990,6 +4990,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,

>      {

>        if (error_operand_p (t))

>  	return error_mark_node;

> +      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)


See above about ort.
Declare simd only allows uniform, linear, aligned, simdlen, inbranch and
notinbranch clauses and none of those support array sections.

> +	  && TREE_CODE (t) == FIELD_DECL)

> +	t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);


handle_omp_array_sections_1 already has recent:
      if (TREE_CODE (t) == FIELD_DECL
          && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY
              || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND))
        ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
so shouldn't that be extended to map/to/from clauses too?
And I guess we should check reduction/in_reduction/task_reduction clauses
too.

> @@ -9003,6 +9037,493 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)

>    return add_stmt (stmt);

>  }

>  

> +/* Used to walk OpenMP target directive body.  */

> +

> +struct omp_target_walk_data

> +{

> +  tree current_object;

> +  bool this_expr_accessed;

> +

> +  hash_map<tree, tree> ptr_members_accessed;

> +  hash_set<tree> lambda_objects_accessed;

> +

> +  tree current_closure;

> +  hash_set<tree> closure_vars_accessed;

> +

> +  hash_set<tree> local_decls;

> +};

> +

> +static tree

> +finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)

> +{

> +  tree t = *tp;

> +  struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;

> +  tree current_object = data->current_object;

> +  tree current_closure = data->current_closure;


This is something that we'll eventually need to do e.g. for declare mapper
in all the 3 FEs, gather what variables might need to be mapped and
for all their types look up the mappers (recursively for nested types and
for all types mentioned in those declare mappers etc.) and remember that
somehow until gimplification.

If it is only preliminary and covers might appear rather than appears,
I think it should be fine.  What this routine does is ultimate, if you see
this somewhere, you say it is accessed, if you see a lambda, again, it has
to be accessed etc.  I'm afraid that is unsafe though.
The IL at this point isn't folded yet, one could have sizeof (this) or other
unevaluated context appear there, or something could appear in a private clause
on some inner construct that doesn't imply an access on the outer target, etc.
So, I think either this function would need to be more careful, especially
for nested OpenMP constructs, or can't it be done through langhooks at
gimplification time when we should know exactly what appears and what
doesn't in the body?

> +  if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t)))


Formatting, missing space before (.

> +	  for (hash_set<tree>::iterator i = data.closure_vars_accessed.begin ();

> +	       i != data.closure_vars_accessed.end (); ++i)

> +	    {

> +	      tree orig_decl = *i;

> +	      tree closure_expr = DECL_VALUE_EXPR (orig_decl);

> +

> +	      if (TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE)

> +		{

> +		  /* this-pointer is processed outside this loop.  */

> +		  if (operand_equal_p (closure_expr, omp_target_this_expr))

> +		    continue;

> +

> +		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);

> +		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);

> +		  OMP_CLAUSE_DECL (c)

> +		    = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR);

> +		  OMP_CLAUSE_SIZE (c) = size_zero_node;

> +		  OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;

> +		  new_clauses.safe_push (c);

> +

> +		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);

> +		  OMP_CLAUSE_SET_MAP_KIND

> +		    (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);

> +		  OMP_CLAUSE_DECL (c) = closure_expr;

> +		  OMP_CLAUSE_SIZE (c) = size_zero_node;

> +		  new_clauses.safe_push (c);

> +		}

> +	      else if (TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE)

> +		{

> +		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);

> +		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);

> +		  OMP_CLAUSE_DECL (c)

> +		    = build1 (INDIRECT_REF,

> +			      TREE_TYPE (TREE_TYPE (closure_expr)),

> +			      closure_expr);

> +		  OMP_CLAUSE_SIZE (c)

> +		    = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr)));

> +		  new_clauses.safe_push (c);

> +

> +		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);

> +		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);

> +		  OMP_CLAUSE_DECL (c) = closure_expr;

> +		  OMP_CLAUSE_SIZE (c) = size_zero_node;

> +		  new_clauses.safe_push (c);


Is it guaranteed everything added here can't have an explicit map clause
already?

	Jakub

Patch

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b90710cba2f..f1f7c43ad30 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15851,7 +15851,8 @@  c_parser_omp_clause_map (c_parser *parser, tree list)
       c_parser_consume_token (parser);
     }
 
-  nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
+  nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
+				   true);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 845d50f3310..f64958eb7ba 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13104,6 +13104,12 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+		  && TREE_CODE (t) == MEM_REF)
+		{
+		  t = TREE_OPERAND (t, 0);
+		  STRIP_NOPS (t);
+		}
 	      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
 		{
 		  if (maybe_ne (mem_ref_offset (t), 0))
@@ -13952,6 +13958,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   tree ordered_clause = NULL_TREE;
   tree schedule_clause = NULL_TREE;
   bool oacc_async = false;
+  bool indir_component_ref_p = false;
   tree last_iterators = NULL_TREE;
   bool last_iterators_remove = false;
   tree *nogroup_seen = NULL;
@@ -14736,6 +14743,11 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    {
 		      while (TREE_CODE (t) == COMPONENT_REF)
 			t = TREE_OPERAND (t, 0);
+		      if (TREE_CODE (t) == MEM_REF)
+			{
+			  t = TREE_OPERAND (t, 0);
+			  STRIP_NOPS (t);
+			}
 		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			  && OMP_CLAUSE_MAP_IMPLICIT (c)
 			  && (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -14802,6 +14814,15 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	       bias) to zero here, so it is not set erroneously to the pointer
 	       size later on in gimplify.c.  */
 	    OMP_CLAUSE_SIZE (c) = size_zero_node;
+	  indir_component_ref_p = false;
+	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	      && TREE_CODE (t) == COMPONENT_REF
+	      && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+	    {
+	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+	      indir_component_ref_p = true;
+	      STRIP_NOPS (t);
+	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
@@ -14874,6 +14895,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 		    || (OMP_CLAUSE_MAP_KIND (c)
 			!= GOMP_MAP_FIRSTPRIVATE_POINTER))
+		   && !indir_component_ref_p
 		   && !c_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -14932,8 +14954,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		}
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
-		   && (ort != C_ORT_OMP
-		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+		   && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 36f99ccf189..6fa7fc26c16 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7597,6 +7597,8 @@  extern tree start_lambda_function		(tree fn, tree lambda_expr);
 extern void finish_lambda_function		(tree body);
 extern bool regenerated_lambda_fn_p		(tree);
 extern tree most_general_lambda			(tree);
+extern tree finish_omp_target			(location_t, tree, tree, bool);
+extern void finish_omp_target_clauses		(location_t, tree, tree *);
 
 /* in tree.c */
 extern int cp_tree_operand_length		(const_tree);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d57ddc4560d..9076bb2fb7e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38201,7 +38201,7 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list)
     }
 
   nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
-					  NULL);
+					  NULL, true);
 
   for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -42225,8 +42225,6 @@  static bool
 cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 		      enum pragma_context context, bool *if_p)
 {
-  tree *pc = NULL, stmt;
-
   if (flag_openmp)
     omp_requires_mask
       = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
@@ -42330,16 +42328,10 @@  cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 		    cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
 		  }
 	    }
-	  tree stmt = make_node (OMP_TARGET);
-	  TREE_TYPE (stmt) = void_type_node;
-	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
-	  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
-	  OMP_TARGET_BODY (stmt) = body;
-	  OMP_TARGET_COMBINED (stmt) = 1;
-	  SET_EXPR_LOCATION (stmt, pragma_tok->location);
-	  add_stmt (stmt);
-	  pc = &OMP_TARGET_CLAUSES (stmt);
-	  goto check_clauses;
+	  c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true);
+	  finish_omp_target (pragma_tok->location,
+			     cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
+	  return true;
 	}
       else if (!flag_openmp)  /* flag_openmp_simd  */
 	{
@@ -42376,49 +42368,13 @@  cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
       return false;
     }
 
-  stmt = make_node (OMP_TARGET);
-  TREE_TYPE (stmt) = void_type_node;
-
-  OMP_TARGET_CLAUSES (stmt)
-    = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
-				 "#pragma omp target", pragma_tok);
-  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
-
-  pc = &OMP_TARGET_CLAUSES (stmt);
+  tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+					    "#pragma omp target", pragma_tok);
+  c_omp_adjust_map_clauses (clauses, true);
   keep_next_level (true);
-  OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
-
-  SET_EXPR_LOCATION (stmt, pragma_tok->location);
-  add_stmt (stmt);
+  tree body = cp_parser_omp_structured_block (parser, if_p);
 
-check_clauses:
-  while (*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_FIRSTPRIVATE_POINTER:
-	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
-	  case GOMP_MAP_ALWAYS_POINTER:
-	  case GOMP_MAP_ATTACH_DETACH:
-	    break;
-	  default:
-	    error_at (OMP_CLAUSE_LOCATION (*pc),
-		      "%<#pragma omp target%> with map-type other "
-		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
-		      "on %<map%> clause");
-	    *pc = OMP_CLAUSE_CHAIN (*pc);
-	    continue;
-	  }
-      pc = &OMP_CLAUSE_CHAIN (*pc);
-    }
+  finish_omp_target (pragma_tok->location, clauses, body, false);
   return true;
 }
 
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 15947b2c812..65621fd6b6f 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18892,6 +18892,11 @@  tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       t = copy_node (t);
       OMP_BODY (t) = stmt;
       OMP_CLAUSES (t) = tmp;
+
+      if (TREE_CODE (t) == OMP_TARGET)
+	finish_omp_target_clauses (EXPR_LOCATION (t), OMP_BODY (t),
+				   &OMP_CLAUSES (t));
+
       if (TREE_CODE (t) == OMP_TARGET && OMP_TARGET_COMBINED (t))
 	{
 	  tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 384c54bd025..455620c0663 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4990,6 +4990,9 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
     {
       if (error_operand_p (t))
 	return error_mark_node;
+      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	  && TREE_CODE (t) == FIELD_DECL)
+	t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
       if (REFERENCE_REF_P (t)
 	  && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	t = TREE_OPERAND (t, 0);
@@ -5018,8 +5021,12 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
-	      if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
-		t = TREE_OPERAND (t, 0);
+	      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+		  && TREE_CODE (t) == INDIRECT_REF)
+		{
+		  t = TREE_OPERAND (t, 0);
+		  STRIP_NOPS (t);
+		}
 	    }
 	  if (REFERENCE_REF_P (t))
 	    t = TREE_OPERAND (t, 0);
@@ -5043,6 +5050,9 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       else if (ort == C_ORT_OMP
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM
 	       && TREE_CODE (t) == PARM_DECL
 	       && DECL_ARTIFICIAL (t)
 	       && DECL_NAME (t) == this_identifier
@@ -5567,6 +5577,8 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	    }
 	  OMP_CLAUSE_DECL (c) = first;
 	  OMP_CLAUSE_SIZE (c) = size;
+	  if (TREE_CODE (t) == FIELD_DECL)
+	    t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
 	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 	      || (TREE_CODE (t) == COMPONENT_REF
 		  && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
@@ -6583,6 +6595,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bool order_seen = false;
   bool schedule_seen = false;
   bool oacc_async = false;
+  bool indir_component_ref_p = false;
   tree last_iterators = NULL_TREE;
   bool last_iterators_remove = false;
   /* 1 if normal/task reduction has been seen, -1 if inscan reduction
@@ -7723,6 +7736,11 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			t = TREE_OPERAND (t, 0);
 		      if (REFERENCE_REF_P (t))
 			t = TREE_OPERAND (t, 0);
+		      if (TREE_CODE (t) == INDIRECT_REF)
+			{
+			  t = TREE_OPERAND (t, 0);
+			  STRIP_NOPS (t);
+			}
 		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			  && OMP_CLAUSE_MAP_IMPLICIT (c)
 			  && (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -7795,10 +7813,15 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      t = TREE_OPERAND (t, 0);
 	      OMP_CLAUSE_DECL (c) = t;
 	    }
+	  indir_component_ref_p = false;
 	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
 	      && TREE_CODE (t) == COMPONENT_REF
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
-	    t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+	    {
+	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+	      indir_component_ref_p = true;
+	      STRIP_NOPS (t);
+	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
 		  || ort == C_ORT_ACC)
@@ -7847,6 +7870,14 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    goto handle_map_references;
 		}
 	    }
+	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	      && !processing_template_decl
+	      && TREE_CODE (t) == FIELD_DECL)
+	    {
+	      OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE,
+								   NULL_TREE);
+	      break;
+	    }
 	  if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
 	    {
 	      if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
@@ -7873,7 +7904,9 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (ort != C_ORT_ACC && t == current_class_ptr)
+	  else if (ort != C_ORT_ACC
+		   && ort != C_ORT_OMP
+		   && t == current_class_ptr)
 	    {
 	      error_at (OMP_CLAUSE_LOCATION (c),
 			"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7886,6 +7919,7 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		   && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 		       || (OMP_CLAUSE_MAP_KIND (c)
 			   != GOMP_MAP_FIRSTPRIVATE_POINTER))
+		   && !indir_component_ref_p
 		   && !cxx_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -9003,6 +9037,493 @@  finish_omp_construct (enum tree_code code, tree body, tree clauses)
   return add_stmt (stmt);
 }
 
+/* Used to walk OpenMP target directive body.  */
+
+struct omp_target_walk_data
+{
+  tree current_object;
+  bool this_expr_accessed;
+
+  hash_map<tree, tree> ptr_members_accessed;
+  hash_set<tree> lambda_objects_accessed;
+
+  tree current_closure;
+  hash_set<tree> closure_vars_accessed;
+
+  hash_set<tree> local_decls;
+};
+
+static tree
+finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
+{
+  tree t = *tp;
+  struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;
+  tree current_object = data->current_object;
+  tree current_closure = data->current_closure;
+
+  if (current_object)
+    {
+      tree this_expr = TREE_OPERAND (current_object, 0);
+
+      if (operand_equal_p (t, this_expr))
+	{
+	  data->this_expr_accessed = true;
+	  *walk_subtrees = 0;
+	  return NULL_TREE;
+	}
+
+      if (TREE_CODE (t) == COMPONENT_REF
+	  && POINTER_TYPE_P (TREE_TYPE (t))
+	  && operand_equal_p (TREE_OPERAND (t, 0), current_object)
+	  && TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL)
+	{
+	  data->this_expr_accessed = true;
+	  tree fld = TREE_OPERAND (t, 1);
+	  if (data->ptr_members_accessed.get (fld) == NULL)
+	    {
+	      if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
+		t = convert_from_reference (t);
+	      data->ptr_members_accessed.put (fld, t);
+	    }
+	  *walk_subtrees = 0;
+	  return NULL_TREE;
+	}
+    }
+
+  /* When the current_function_decl is a lambda function, the closure object
+     argument's type seems to not yet have fields layed out, so a recording
+     of DECL_VALUE_EXPRs during the target body walk seems the only way to
+     find them.  */
+  if (current_closure
+      && (TREE_CODE (t) == VAR_DECL
+	  || TREE_CODE (t) == PARM_DECL
+	  || TREE_CODE (t) == RESULT_DECL)
+      && DECL_HAS_VALUE_EXPR_P (t)
+      && TREE_CODE (DECL_VALUE_EXPR (t)) == COMPONENT_REF
+      && operand_equal_p (current_closure,
+			  TREE_OPERAND (DECL_VALUE_EXPR (t), 0)))
+    {
+      if (!data->closure_vars_accessed.contains (t))
+	data->closure_vars_accessed.add (t);
+      *walk_subtrees = 0;
+      return NULL_TREE;
+    }
+
+  if (TREE_CODE (t) == BIND_EXPR)
+    {
+      tree block = BIND_EXPR_BLOCK (t);
+      for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var))
+	if (!data->local_decls.contains (var))
+	  data->local_decls.add (var);
+      return NULL_TREE;
+    }
+
+  if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t)))
+    {
+      tree lt = TREE_TYPE (t);
+      gcc_assert (CLASS_TYPE_P (lt));
+
+      if (!data->lambda_objects_accessed.contains (t)
+	  /* Do not prepare to create target maps for locally declared
+	     lambdas or anonymous ones.  */
+	  && !data->local_decls.contains (t)
+	  && TREE_CODE (t) != TARGET_EXPR)
+	data->lambda_objects_accessed.add (t);
+      *walk_subtrees = 0;
+      return NULL_TREE;
+    }
+
+  return NULL_TREE;
+}
+
+void
+finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
+{
+  omp_target_walk_data data;
+  data.this_expr_accessed = false;
+
+  tree ct = current_nonlambda_class_type ();
+  if (ct)
+    {
+      tree object = maybe_dummy_object (ct, NULL);
+      object = maybe_resolve_dummy (object, true);
+      data.current_object = object;
+    }
+  else
+    data.current_object = NULL_TREE;
+
+  if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
+    {
+      tree closure = DECL_ARGUMENTS (current_function_decl);
+      data.current_closure = build_indirect_ref (loc, closure, RO_UNARY_STAR);
+    }
+  else
+    data.current_closure = NULL_TREE;
+
+  cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
+
+  auto_vec<tree, 16> new_clauses;
+
+  if (data.this_expr_accessed)
+    {
+      tree omp_target_this_expr = TREE_OPERAND (data.current_object, 0);
+
+      /* See if explicit user-specified map(this[:]) clause already exists.
+	 If not, we create an implicit map(tofrom:this[:1]) clause.  */
+      tree *explicit_this_deref_map = NULL;
+      for (tree *c = clauses_ptr; *c; c = &OMP_CLAUSE_CHAIN (*c))
+	if (OMP_CLAUSE_CODE (*c) == OMP_CLAUSE_MAP
+	    && TREE_CODE (OMP_CLAUSE_DECL (*c)) == INDIRECT_REF
+	    && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*c), 0),
+				omp_target_this_expr))
+	  {
+	    explicit_this_deref_map = c;
+	    break;
+	  }
+
+      if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
+	{
+	  /* For lambda functions, we need to first create a copy of the
+	     __closure object.  */
+	  tree closure = DECL_ARGUMENTS (current_function_decl);
+	  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	  OMP_CLAUSE_DECL (c)
+	    = build_indirect_ref (loc, closure, RO_UNARY_STAR);
+	  OMP_CLAUSE_SIZE (c)
+	    = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure)));
+	  new_clauses.safe_push (c);
+
+	  tree closure_obj = OMP_CLAUSE_DECL (c);
+	  tree closure_type = TREE_TYPE (closure_obj);
+
+	  gcc_assert (LAMBDA_TYPE_P (closure_type)
+		      && CLASS_TYPE_P (closure_type));
+
+	  tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  OMP_CLAUSE_DECL (c2) = closure;
+	  OMP_CLAUSE_SIZE (c2) = size_zero_node;
+	  new_clauses.safe_push (c2);
+
+	  STRIP_NOPS (omp_target_this_expr);
+	  gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr));
+	  omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr);
+
+	  for (hash_set<tree>::iterator i = data.closure_vars_accessed.begin ();
+	       i != data.closure_vars_accessed.end (); ++i)
+	    {
+	      tree orig_decl = *i;
+	      tree closure_expr = DECL_VALUE_EXPR (orig_decl);
+
+	      if (TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE)
+		{
+		  /* this-pointer is processed outside this loop.  */
+		  if (operand_equal_p (closure_expr, omp_target_this_expr))
+		    continue;
+
+		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+		  OMP_CLAUSE_DECL (c)
+		    = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR);
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+		  new_clauses.safe_push (c);
+
+		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND
+		    (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+		  OMP_CLAUSE_DECL (c) = closure_expr;
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  new_clauses.safe_push (c);
+		}
+	      else if (TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE)
+		{
+		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+		  OMP_CLAUSE_DECL (c)
+		    = build1 (INDIRECT_REF,
+			      TREE_TYPE (TREE_TYPE (closure_expr)),
+			      closure_expr);
+		  OMP_CLAUSE_SIZE (c)
+		    = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr)));
+		  new_clauses.safe_push (c);
+
+		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
+		  OMP_CLAUSE_DECL (c) = closure_expr;
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  new_clauses.safe_push (c);
+		}
+	    }
+
+	  if (explicit_this_deref_map)
+	    {
+	      /* Transform *this into *__closure->this in maps.  */
+	      tree this_map = *explicit_this_deref_map;
+	      OMP_CLAUSE_DECL (this_map)
+		= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
+
+	      tree nc = OMP_CLAUSE_CHAIN (this_map);
+	      gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+			  && (OMP_CLAUSE_MAP_KIND (nc)
+			      == GOMP_MAP_FIRSTPRIVATE_POINTER));
+	      OMP_CLAUSE_DECL (nc) = omp_target_this_expr;
+	      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_POINTER);
+
+	      /* Unlink this two-map sequence away from the chain.  */
+	      *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc);
+
+	      /* Move map(*__closure->this) map(always_pointer:__closure->this)
+		 sequence to right after __closure map.  */
+	      new_clauses.safe_push (this_map);
+	      new_clauses.safe_push (nc);
+	    }
+	  else
+	    {
+	      tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM);
+	      OMP_CLAUSE_DECL (c3)
+		= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
+	      OMP_CLAUSE_SIZE (c3)
+		= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+
+	      tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
+
+	      OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
+	      OMP_CLAUSE_SIZE (c4) = size_zero_node;
+
+	      new_clauses.safe_push (c3);
+	      new_clauses.safe_push (c4);
+	    }
+	}
+      else
+	{
+	  /* For the non-lambda case, we only need to create map(this[:1]) when
+	     it's not present, no transforming needed.  */
+	  if (!explicit_this_deref_map)
+	    {
+	      tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+	      OMP_CLAUSE_DECL (c)
+		= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
+	      OMP_CLAUSE_SIZE (c)
+		= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+
+	      tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	      STRIP_NOPS (omp_target_this_expr);
+	      OMP_CLAUSE_DECL (c2) = omp_target_this_expr;
+	      OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+	      new_clauses.safe_push (c);
+	      new_clauses.safe_push (c2);
+	    }
+	}
+
+      if (!data.ptr_members_accessed.is_empty ())
+	for (hash_map<tree, tree>::iterator i
+	       = data.ptr_members_accessed.begin ();
+	     i != data.ptr_members_accessed.end (); ++i)
+	  {
+	    /* For each referenced member that is of pointer or
+	       reference-to-pointer type, create the equivalent of
+	       map(alloc:this->ptr[:0]).  */
+	    tree field_decl = (*i).first;
+	    tree ptr_member = (*i).second;
+
+	    for (tree c = *clauses_ptr; c; c = OMP_CLAUSE_CHAIN (c))
+	      {
+		/* If map(this->ptr[:N] already exists, avoid creating another
+		   such map.  */
+		tree decl = OMP_CLAUSE_DECL (c);
+		if ((TREE_CODE (decl) == INDIRECT_REF
+		     || TREE_CODE (decl) == MEM_REF)
+		    && operand_equal_p (TREE_OPERAND (decl, 0),
+					ptr_member))
+		  goto next_ptr_member;
+	      }
+
+	    if (!cxx_mark_addressable (ptr_member))
+	      gcc_unreachable ();
+
+	    if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE)
+	      {
+		/* For reference to pointers, we need to map the referenced
+		   pointer first for things to be correct.  */
+		tree ptr_member_type = TREE_TYPE (ptr_member);
+
+		/* Map pointer target as zero-length array section.  */
+		tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+		OMP_CLAUSE_DECL (c)
+		  = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member);
+		OMP_CLAUSE_SIZE (c) = size_zero_node;
+		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+		/* Map pointer to zero-length array section.  */
+		tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		OMP_CLAUSE_SET_MAP_KIND
+		  (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION);
+		OMP_CLAUSE_DECL (c2) = ptr_member;
+		OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+		/* Attach reference-to-pointer field to pointer.  */
+		tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH);
+		OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0);
+		OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+		new_clauses.safe_push (c);
+		new_clauses.safe_push (c2);
+		new_clauses.safe_push (c3);
+	      }
+	    else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE)
+	      {
+		/* Map pointer target as zero-length array section.  */
+		tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+		OMP_CLAUSE_DECL (c)
+		  = build_indirect_ref (loc, ptr_member, RO_UNARY_STAR);
+		OMP_CLAUSE_SIZE (c) = size_zero_node;
+		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+		/* Attach zero-length array section to pointer.  */
+		tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		OMP_CLAUSE_SET_MAP_KIND
+		  (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+		OMP_CLAUSE_DECL (c2) = ptr_member;
+		OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+		new_clauses.safe_push (c);
+		new_clauses.safe_push (c2);
+	      }
+	    else
+	      gcc_unreachable ();
+
+	  next_ptr_member:
+	    ;
+	  }
+    }
+
+  if (!data.lambda_objects_accessed.is_empty ())
+    {
+      for (hash_set<tree>::iterator i = data.lambda_objects_accessed.begin ();
+	   i != data.lambda_objects_accessed.end (); ++i)
+	{
+	  tree lobj = *i;
+	  if (TREE_CODE (lobj) == TARGET_EXPR)
+	    lobj = TREE_OPERAND (lobj, 0);
+
+	  tree lt = TREE_TYPE (lobj);
+	  gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt));
+
+	  tree lc = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (lc, GOMP_MAP_TO);
+	  OMP_CLAUSE_DECL (lc) = lobj;
+	  OMP_CLAUSE_SIZE (lc) = TYPE_SIZE_UNIT (lt);
+	  new_clauses.truncate (0);
+	  new_clauses.safe_push (lc);
+
+	  for (tree fld = TYPE_FIELDS (lt); fld; fld = DECL_CHAIN (fld))
+	    {
+	      if (TREE_CODE (TREE_TYPE (fld)) == POINTER_TYPE)
+		{
+		  tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
+				     lobj, fld, NULL_TREE);
+		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+		  OMP_CLAUSE_DECL (c)
+		    = build_indirect_ref (loc, exp, RO_UNARY_STAR);
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+		  new_clauses.safe_push (c);
+
+		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND
+		    (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+		  OMP_CLAUSE_DECL (c) = exp;
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  new_clauses.safe_push (c);
+		}
+	      else if (TREE_CODE (TREE_TYPE (fld)) == REFERENCE_TYPE)
+		{
+		  tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
+				     lobj, fld, NULL_TREE);
+		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+		  OMP_CLAUSE_DECL (c)
+		    = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp);
+		  OMP_CLAUSE_SIZE (c)
+		    = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (exp)));
+		  new_clauses.safe_push (c);
+
+		  c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
+		  OMP_CLAUSE_DECL (c) = exp;
+		  OMP_CLAUSE_SIZE (c) = size_zero_node;
+		  new_clauses.safe_push (c);
+		}
+	    }
+	}
+    }
+
+  tree c = *clauses_ptr;
+  for (int i = new_clauses.length () - 1; i >= 0; i--)
+    {
+      OMP_CLAUSE_CHAIN (new_clauses[i]) = c;
+      c = new_clauses[i];
+    }
+  *clauses_ptr = c;
+}
+
+tree
+finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
+{
+  if (!processing_template_decl)
+    finish_omp_target_clauses (loc, body, &clauses);
+
+  tree stmt = make_node (OMP_TARGET);
+  TREE_TYPE (stmt) = void_type_node;
+  OMP_TARGET_CLAUSES (stmt) = clauses;
+  OMP_TARGET_BODY (stmt) = body;
+  OMP_TARGET_COMBINED (stmt) = combined_p;
+  SET_EXPR_LOCATION (stmt, loc);
+
+  tree c = clauses;
+  while (c)
+    {
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	switch (OMP_CLAUSE_MAP_KIND (c))
+	  {
+	  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_FIRSTPRIVATE_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	    break;
+	  default:
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "%<#pragma omp target%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    break;
+	  }
+      c = OMP_CLAUSE_CHAIN (c);
+    }
+  return add_stmt (stmt);
+}
+
 tree
 finish_omp_parallel (tree clauses, tree body)
 {
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2606998152d..86d328c99fc 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -53,6 +53,7 @@  along with GCC; see the file COPYING3.  If not see
 #include "langhooks.h"
 #include "tree-cfg.h"
 #include "tree-ssa.h"
+#include "tree-hash-traits.h"
 #include "omp-general.h"
 #include "omp-low.h"
 #include "gimple-low.h"
@@ -8738,7 +8739,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 {
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
-  hash_map<tree, tree> *struct_map_to_clause = NULL;
+  hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
   hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL, *orig_list_p = list_p;
   int handled_depend_iterators = -1;
@@ -9176,7 +9177,14 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
 
-	  if (!DECL_P (decl))
+	  if (TREE_CODE (decl) == TARGET_EXPR)
+	    {
+	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+				 is_gimple_lvalue, fb_lvalue)
+		  == GS_ERROR)
+		remove = true;
+	    }
+	  else if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
 	      if (TREE_CODE (d) == ARRAY_REF)
@@ -9192,12 +9200,15 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  && TREE_CODE (decl) == INDIRECT_REF
 		  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
 		  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
-		      == REFERENCE_TYPE))
+		      == REFERENCE_TYPE)
+		  && (OMP_CLAUSE_MAP_KIND (c)
+		      != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
 		{
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
 	      bool indir_p = false;
+	      bool component_ref_p = false;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
 	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9208,6 +9219,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    {
 		      decl = TREE_OPERAND (decl, 0);
+		      component_ref_p = true;
 		      if (((TREE_CODE (decl) == MEM_REF
 			    && integer_zerop (TREE_OPERAND (decl, 1)))
 			   || INDIRECT_REF_P (decl))
@@ -9216,6 +9228,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			{
 			  indir_p = true;
 			  decl = TREE_OPERAND (decl, 0);
+			  STRIP_NOPS (decl);
 			}
 		      if (TREE_CODE (decl) == INDIRECT_REF
 			  && DECL_P (TREE_OPERAND (decl, 0))
@@ -9227,8 +9240,11 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		    }
 		}
-	      else if (TREE_CODE (decl) == COMPONENT_REF)
+	      else if (TREE_CODE (decl) == COMPONENT_REF
+		       && (OMP_CLAUSE_MAP_KIND (c)
+			   != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
 		{
+		  component_ref_p = true;
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    decl = TREE_OPERAND (decl, 0);
 		  if (TREE_CODE (decl) == INDIRECT_REF
@@ -9298,7 +9314,10 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      if (code == OACC_UPDATE
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
-	      if (DECL_P (decl)
+	      if ((DECL_P (decl)
+		   || (component_ref_p
+		       && (INDIRECT_REF_P (decl)
+			   || TREE_CODE (decl) == MEM_REF)))
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9355,7 +9374,10 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  gcc_assert (base == decl);
 
 		  splay_tree_node n
-		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+		    = (DECL_P (decl)
+		       ? splay_tree_lookup (ctx->variables,
+					    (splay_tree_key) decl)
+		       : NULL);
 		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)
 			      == GOMP_MAP_ALWAYS_POINTER);
 		  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
@@ -9381,7 +9403,11 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		      has_attachments = true;
 		    }
-		  if (n == NULL || (n->value & GOVD_MAP) == 0)
+		  if ((DECL_P (decl)
+		       && (n == NULL || (n->value & GOVD_MAP) == 0))
+		      || (!DECL_P (decl)
+			  && (!struct_map_to_clause
+			      || struct_map_to_clause->get (decl) == NULL)))
 		    {
 		      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 						 OMP_CLAUSE_MAP);
@@ -9392,7 +9418,18 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      if (base_ref)
 			OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
 		      else
-			OMP_CLAUSE_DECL (l) = decl;
+			{
+			  OMP_CLAUSE_DECL (l) = unshare_expr (decl);
+			  if (!DECL_P (OMP_CLAUSE_DECL (l))
+			      && (gimplify_expr (&OMP_CLAUSE_DECL (l),
+						 pre_p, NULL, is_gimple_lvalue,
+						 fb_lvalue)
+				  == GS_ERROR))
+			    {
+			      remove = true;
+			      break;
+			    }
+			}
 		      OMP_CLAUSE_SIZE (l)
 			= (!attach
 			   ? size_int (1)
@@ -9400,7 +9437,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
 			   : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
 		      if (struct_map_to_clause == NULL)
-			struct_map_to_clause = new hash_map<tree, tree>;
+			struct_map_to_clause
+			  = new hash_map<tree_operand_hash, tree>;
 		      struct_map_to_clause->put (decl, l);
 		      if (ptr || attach_detach)
 			{
@@ -9434,7 +9472,32 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			flags |= GOVD_SEEN;
 		      if (has_attachments)
 			flags |= GOVD_MAP_HAS_ATTACHMENTS;
-		      goto do_add_decl;
+
+		      /* If this is a *pointer-to-struct expression, make sure a
+			 firstprivate map of the base-pointer exists.  */
+		      if (component_ref_p
+			  && ((TREE_CODE (decl) == MEM_REF
+			       && integer_zerop (TREE_OPERAND (decl, 1)))
+			      || INDIRECT_REF_P (decl))
+			  && DECL_P (TREE_OPERAND (decl, 0))
+			  && !splay_tree_lookup (ctx->variables,
+						 ((splay_tree_key)
+						  TREE_OPERAND (decl, 0))))
+			{
+			  decl = TREE_OPERAND (decl, 0);
+			  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						      OMP_CLAUSE_MAP);
+			  enum gomp_map_kind mkind
+			    = GOMP_MAP_FIRSTPRIVATE_POINTER;
+			  OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+			  OMP_CLAUSE_DECL (c2) = decl;
+			  OMP_CLAUSE_SIZE (c2) = size_zero_node;
+			  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = c2;
+			}
+
+		      if (DECL_P (decl))
+			goto do_add_decl;
 		    }
 		  else if (struct_map_to_clause)
 		    {
@@ -9543,6 +9606,13 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		      else if (*sc != c)
 			{
+			  if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+					     fb_lvalue)
+			      == GS_ERROR)
+			    {
+			      remove = true;
+			      break;
+			    }
 			  *list_p = OMP_CLAUSE_CHAIN (c);
 			  OMP_CLAUSE_CHAIN (c) = *sc;
 			  *sc = c;
@@ -9570,6 +9640,24 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  break;
 		}
 
+	      /* If this was of the form map(*pointer_to_struct), then the
+		 'pointer_to_struct' DECL should be considered deref'ed.  */
+	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+		   || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c))
+		   || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c)))
+		  && INDIRECT_REF_P (orig_decl)
+		  && DECL_P (TREE_OPERAND (orig_decl, 0))
+		  && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE)
+		{
+		  tree ptr = TREE_OPERAND (orig_decl, 0);
+		  if (!struct_deref_set || !struct_deref_set->contains (ptr))
+		    {
+		      if (!struct_deref_set)
+			struct_deref_set = new hash_set<tree> ();
+		      struct_deref_set->add (ptr);
+		    }
+		}
+
 	      if (!remove
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -10846,6 +10934,12 @@  gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		    }
 		}
 	    }
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+	      && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA))
+	    {
+	      remove = true;
+	      break;
+	    }
 	  if (!DECL_P (decl))
 	    {
 	      if ((ctx->region_type & ORT_TARGET) != 0
@@ -10892,10 +10986,6 @@  gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		      = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
 		}
 	    }
-	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-		   && (code == OMP_TARGET_EXIT_DATA
-		       || code == OACC_EXIT_DATA))
-	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
 		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 2cc2a186080..0307e97512b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12268,6 +12268,8 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH:
 	  case GOMP_MAP_DETACH:
+	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
 	    break;
 	  case GOMP_MAP_IF_PRESENT:
 	  case GOMP_MAP_FORCE_ALLOC:
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C
new file mode 100644
index 00000000000..f4d40ec8e4b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -0,0 +1,36 @@ 
+// { dg-do compile }
+// { dg-options "-fopenmp -fdump-tree-gimple" }
+
+struct S
+{
+  int a, b;
+  void bar (int);
+};
+
+void
+S::bar (int x)
+{
+  #pragma omp target map (alloc: a, b)
+    ;
+  #pragma omp target enter data map (alloc: a, b)
+}
+
+template <int N>
+struct T
+{
+  int a, b;
+  void bar (int);
+};
+
+template <int N>
+void
+T<N>::bar (int x)
+{
+  #pragma omp target map (alloc: a, b)
+    ;
+  #pragma omp target enter data map (alloc: a, b)
+}
+
+template struct T<0>;
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
new file mode 100644
index 00000000000..7dceef80f47
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -0,0 +1,94 @@ 
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+
+template <typename L>
+void
+omp_target_loop (int begin, int end, L loop)
+{
+  #pragma omp target teams distribute parallel for
+  for (int i = begin; i < end; i++)
+    loop (i);
+}
+
+struct S
+{
+  int a, len;
+  int *ptr;
+
+  auto merge_data_func (int *iptr, int &b)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  mapped = (ptr != NULL && iptr != NULL);
+	  if (mapped)
+	    {
+	      for (int i = 0; i < len; i++)
+		ptr[i] += a + b + iptr[i];
+	    }
+	}
+	return mapped;
+      };
+    return fn;
+  }
+};
+
+int x = 1;
+
+int main (void)
+{
+  const int N = 10;
+  int *data1 = new int[N];
+  int *data2 = new int[N];
+  memset (data1, 0xab, sizeof (int) * N);
+  memset (data1, 0xcd, sizeof (int) * N);
+
+  int val = 1;
+  int &valref = val;
+  #pragma omp target enter data map(alloc: data1[:N], data2[:N])
+
+  omp_target_loop (0, N, [=](int i) { data1[i] = val; });
+  omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
+
+  #pragma omp target update from(data1[:N], data2[:N])
+
+  for (int i = 0; i < N; i++)
+    {
+      if (data1[i] != 1) abort ();
+      if (data2[i] != 2) abort ();
+    }
+
+  #pragma omp target exit data map(delete: data1[:N], data2[:N])
+
+  int b = 8;
+  S s = { 4, N, data1 };
+  auto f = s.merge_data_func (data2, b);
+
+  if (f ()) abort ();
+
+  #pragma omp target enter data map(to: data1[:N])
+  if (f ()) abort ();
+
+  #pragma omp target enter data map(to: data2[:N])
+  if (!f ()) abort ();
+
+  #pragma omp target exit data map(from: data1[:N], data2[:N])
+
+  for (int i = 0; i < N; i++)
+    {
+      if (data1[i] != 0xf) abort ();
+      if (data2[i] != 2) abort ();
+    }
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
new file mode 100644
index 00000000000..bdf2564cd04
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
@@ -0,0 +1,35 @@ 
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}
+
+/* Make sure lambda objects do NOT appear in target maps.  */
+/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C
new file mode 100644
index 00000000000..de93a3e5e57
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C
@@ -0,0 +1,33 @@ 
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+struct S
+{
+  int a, b, c, d;
+
+  int sum (void)
+  {
+    int val = 0;
+    val += a + b + this->c + this->d;
+    return val;
+  }
+
+  int sum_offload (void)
+  {
+    int val = 0;
+    #pragma omp target map(val)
+    val += a + b + this->c + this->d;
+    return val;
+  }
+};
+
+int main (void)
+{
+  S s = { 1, 2, 3, 4 };
+  if (s.sum () != s.sum_offload ())
+    abort ();
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
new file mode 100644
index 00000000000..679c85a54dd
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -0,0 +1,49 @@ 
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-do compile }
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+
+extern "C" void abort ();
+
+struct T
+{
+  int x, y;
+
+  auto sum_func (int n)
+  {
+    auto fn = [=](int m) -> int
+      {
+	int v;
+	v = (x + y) * n + m;
+	return v;
+      };
+    return fn;
+  }
+
+  auto sum_func_offload (int n)
+  {
+    auto fn = [=](int m) -> int
+      {
+	int v;
+	#pragma omp target map(from:v)
+	v = (x + y) * n + m;
+	return v;
+      };
+    return fn;
+  }
+
+};
+
+int main (void)
+{
+  T a = { 1, 2 };
+
+  auto s1 = a.sum_func (3);
+  auto s2 = a.sum_func_offload (3);
+
+  if (s1 (1) != s2 (1))
+    abort ();
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
new file mode 100644
index 00000000000..08568f9284c
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -0,0 +1,105 @@ 
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+extern "C" void abort ();
+
+struct S
+{
+  int * ptr;
+  int ptr_len;
+
+  int *&refptr;
+  int refptr_len;
+
+  bool set_ptr (int n)
+  {
+    bool mapped;
+    #pragma omp target map(from:mapped)
+    {
+      if (ptr != NULL)
+	for (int i = 0; i < ptr_len; i++)
+	  ptr[i] = n;
+      mapped = (ptr != NULL);
+    }
+    return mapped;
+  }
+
+  bool set_refptr (int n)
+  {
+    bool mapped;
+    #pragma omp target map(from:mapped)
+    {
+      if (refptr != NULL)
+	for (int i = 0; i < refptr_len; i++)
+	  refptr[i] = n;
+      mapped = (refptr != NULL);
+    }
+    return mapped;
+  }
+};
+
+int main (void)
+{
+  #define N 10
+  int *ptr1 = new int[N];
+  int *ptr2 = new int[N];
+
+  memset (ptr1, 0, sizeof (int) * N);
+  memset (ptr2, 0, sizeof (int) * N);
+
+  S s = { ptr1, N, ptr2, N };
+
+  bool mapped;
+  int val = 123;
+
+  mapped = s.set_ptr (val);
+  if (mapped)
+    abort ();
+  if (s.ptr != ptr1)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 0)
+      abort ();
+
+  mapped = s.set_refptr (val);
+  if (mapped)
+    abort ();
+  if (s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 0)
+      abort ();
+
+  #pragma omp target data map(ptr1[:N])
+  mapped = s.set_ptr (val);
+
+  if (!mapped)
+    abort ();
+  if (s.set_refptr (0))
+    abort ();
+  if (s.ptr != ptr1 || s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != val)
+      abort ();
+
+  #pragma omp target data map(ptr2[:N])
+  mapped = s.set_refptr (val);
+
+  if (!mapped)
+    abort ();
+  if (s.set_ptr (0))
+    abort ();
+  if (s.ptr != ptr1 || s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != val)
+      abort ();
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
new file mode 100644
index 00000000000..3b2d5811350
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -0,0 +1,107 @@ 
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+
+struct T
+{
+  int *ptr;
+  int ptr_len;
+
+  int *&refptr;
+  int refptr_len;
+
+  auto set_ptr_func (int n)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  if (ptr)
+	    for (int i = 0; i < ptr_len; i++)
+	      ptr[i] = n;
+	  mapped = (ptr != NULL);
+	}
+	return mapped;
+      };
+    return fn;
+  }
+
+  auto set_refptr_func (int n)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  if (refptr)
+	    for (int i = 0; i < refptr_len; i++)
+	      refptr[i] = n;
+	  mapped = (refptr != NULL);
+	}
+	return mapped;
+      };
+    return fn;
+  }
+};
+
+int main (void)
+{
+  #define N 10
+  int *ptr1 = new int[N];
+  int *ptr2 = new int[N];
+
+  memset (ptr1, 0, sizeof (int) * N);
+  memset (ptr2, 0, sizeof (int) * N);
+
+  T a = { ptr1, N, ptr2, N };
+
+  auto p1 = a.set_ptr_func (1);
+  auto r2 = a.set_refptr_func (2);
+
+  if (p1 ())
+    abort ();
+  if (r2 ())
+    abort ();
+
+  if (a.ptr != ptr1)
+    abort ();
+  if (a.refptr != ptr2)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 0)
+      abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 0)
+      abort ();
+
+  #pragma omp target data map(ptr1[:N], ptr2[:N])
+  {
+    if (!p1 ())
+      abort ();
+    if (!r2 ())
+      abort ();
+  }
+
+  if (a.ptr != ptr1)
+    abort ();
+  if (a.refptr != ptr2)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 1)
+      abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 2)
+      abort ();
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C
new file mode 100644
index 00000000000..a9ac74bcf1f
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C
@@ -0,0 +1,34 @@ 
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+  T a, b, c, d;
+
+  T sum (void)
+  {
+    T val = 0;
+    val += a + b + this->c + this->d;
+    return val;
+  }
+
+  T sum_offload (void)
+  {
+    T val = 0;
+    #pragma omp target map(val)
+    val += a + b + this->c + this->d;
+    return val;
+  }
+};
+
+int main (void)
+{
+  S<int> s = { 1, 2, 3, 4 };
+  if (s.sum () != s.sum_offload ())
+    abort ();
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C
index d03b8a0728e..b521a4faf5e 100644
--- a/gcc/testsuite/g++.dg/gomp/this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/this-2.C
@@ -9,14 +9,14 @@  struct S
 void
 S::bar (int x)
 {
-  #pragma omp target map (this, x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this, x)		// { dg-error "cannot take the address of .this., which is an rvalue expression" }
     ;
-  #pragma omp target map (this[0], x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this[0], x)
     ;
-  #pragma omp target update to (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update to (this[0], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this[1], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target update to (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update to (this[0], x)
+  #pragma omp target update from (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update from (this[1], x)
 }
 
 template <int N>
@@ -29,14 +29,14 @@  template <int N>
 void
 T<N>::bar (int x)
 {
-  #pragma omp target map (this, x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this, x)		// { dg-error "cannot take the address of .this., which is an rvalue expression" }
     ;
-  #pragma omp target map (this[0], x)		// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target map (this[0], x)
     ;
-  #pragma omp target update to (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update to (this[0], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this, x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
-  #pragma omp target update from (this[1], x)	// { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+  #pragma omp target update to (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update to (this[0], x)
+  #pragma omp target update from (this, x)	// { dg-error "cannot take the address of .this., which is an rvalue expression" }
+  #pragma omp target update from (this[1], x)
 }
 
 template struct T<0>;
diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c
new file mode 100644
index 00000000000..3e7921270c9
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/target-3.c
@@ -0,0 +1,16 @@ 
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+struct S
+{
+  int a, b;
+};
+
+void foo (struct S *s)
+{
+  #pragma omp target map (alloc: s->a, s->b)
+    ;
+  #pragma omp target enter data map (alloc: s->a, s->b)
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index bcbe669653c..6a855baa26e 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -836,6 +836,7 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	{
 	case GOMP_MAP_ALLOC:
 	case GOMP_MAP_POINTER:
+	case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
 	  pp_string (pp, "alloc");
 	  break;
 	case GOMP_MAP_IF_PRESENT:
@@ -914,6 +915,9 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_ATTACH_DETACH:
 	  pp_string (pp, "attach_detach");
 	  break;
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  pp_string (pp, "attach_zero_length_array_section");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -932,6 +936,9 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	    case GOMP_MAP_ALWAYS_POINTER:
 	      pp_string (pp, " [pointer assign, bias: ");
 	      break;
+	    case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	      pp_string (pp, " [pointer assign, zero-length array section, bias: ");
+	      break;
 	    case GOMP_MAP_TO_PSET:
 	      pp_string (pp, " [pointer set, len: ");
 	      break;
@@ -939,6 +946,7 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	    case GOMP_MAP_DETACH:
 	    case GOMP_MAP_FORCE_DETACH:
 	    case GOMP_MAP_ATTACH_DETACH:
+	    case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	      pp_string (pp, " [bias: ");
 	      break;
 	    default:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 6e163b02560..4a4a14393d2 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -132,6 +132,11 @@  enum gomp_map_kind
        No refcount is bumped by this, and the store is done unconditionally.  */
     GOMP_MAP_ALWAYS_POINTER =		(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_FLAG_SPECIAL | 1),
+    /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to
+       NULL if target is not mapped.  */
+    GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
+      =					(GOMP_MAP_FLAG_SPECIAL_2
+					 | GOMP_MAP_FLAG_SPECIAL | 2),
     /* Forced deallocation of zero length array section.  */
     GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
       =					(GOMP_MAP_FLAG_SPECIAL_2
@@ -152,6 +157,12 @@  enum gomp_map_kind
     GOMP_MAP_FORCE_DETACH =		(GOMP_MAP_DEEP_COPY
 					 | GOMP_MAP_FLAG_FORCE | 1),
 
+    /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections
+       (i.e. set to NULL when array section is not mapped) Currently only used
+       by OpenMP.  */
+    GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
+      =					(GOMP_MAP_DEEP_COPY | 2),
+
     /* Internal to GCC, not used in libgomp.  */
     /* Do not map, but pointer assign a pointer instead.  */
     GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1),
@@ -175,7 +186,8 @@  enum gomp_map_kind
   ((X) == GOMP_MAP_ALWAYS_POINTER)
 
 #define GOMP_MAP_POINTER_P(X) \
-  ((X) == GOMP_MAP_POINTER)
+  ((X) == GOMP_MAP_POINTER \
+   || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
 
 #define GOMP_MAP_ALWAYS_TO_P(X) \
   (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 8d25dc8e2a8..235e919f191 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1234,7 +1234,7 @@  extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
 extern void gomp_attach_pointer (struct gomp_device_descr *,
 				 struct goacc_asyncqueue *, splay_tree,
 				 splay_tree_key, uintptr_t, size_t,
-				 struct gomp_coalesce_buf *);
+				 struct gomp_coalesce_buf *, bool);
 extern void gomp_detach_pointer (struct gomp_device_descr *,
 				 struct goacc_asyncqueue *, splay_tree_key,
 				 uintptr_t, bool, struct gomp_coalesce_buf *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index c21508f3739..5ee8647cddf 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -937,7 +937,7 @@  acc_attach_async (void **hostaddr, int async)
     }
 
   gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
-		       0, NULL);
+		       0, NULL, false);
 
   gomp_mutex_unlock (&acc_dev->lock);
 }
@@ -1141,7 +1141,7 @@  goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
 	    {
 	      gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
-				   (uintptr_t) h, s, NULL);
+				   (uintptr_t) h, s, NULL, false);
 	      /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
 		 reference counts ('n->refcount', 'n->dynamic_refcount').  */
 	    }
@@ -1159,7 +1159,8 @@  goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		splay_tree_key m
 		  = lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
 		gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
-				     (uintptr_t) hostaddrs[j], sizes[j], NULL);
+				     (uintptr_t) hostaddrs[j], sizes[j], NULL,
+				     false);
 	      }
 
 	  bool processed = false;
diff --git a/libgomp/target.c b/libgomp/target.c
index bb09d501dd6..1ff6a503fe4 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -496,7 +496,8 @@  gomp_map_vars_existing (struct gomp_device_descr *devicep,
 			struct gomp_coalesce_buf *cbuf,
 			htab_t *refcount_set)
 {
-  assert (kind != GOMP_MAP_ATTACH);
+  assert (kind != GOMP_MAP_ATTACH
+	  || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
 
   tgt_var->key = oldn;
   tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
@@ -536,7 +537,8 @@  get_kind (bool short_mapkind, void *kinds, int idx)
 static void
 gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
 		  uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias,
-		  struct gomp_coalesce_buf *cbuf)
+		  struct gomp_coalesce_buf *cbuf,
+		  bool allow_zero_length_array_sections)
 {
   struct gomp_device_descr *devicep = tgt->device_descr;
   struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -558,16 +560,24 @@  gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
   splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n == NULL)
     {
-      gomp_mutex_unlock (&devicep->lock);
-      gomp_fatal ("Pointer target of array section wasn't mapped");
-    }
-  cur_node.host_start -= n->host_start;
-  cur_node.tgt_offset
-    = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
-  /* At this point tgt_offset is target address of the
-     array section.  Now subtract bias to get what we want
-     to initialize the pointer with.  */
-  cur_node.tgt_offset -= bias;
+      if (allow_zero_length_array_sections)
+	cur_node.tgt_offset = 0;
+      else
+	{
+	  gomp_mutex_unlock (&devicep->lock);
+	  gomp_fatal ("Pointer target of array section wasn't mapped");
+	}
+    }
+  else
+    {
+      cur_node.host_start -= n->host_start;
+      cur_node.tgt_offset
+	= n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
+      /* At this point tgt_offset is target address of the
+	 array section.  Now subtract bias to get what we want
+	 to initialize the pointer with.  */
+      cur_node.tgt_offset -= bias;
+    }
   gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset),
 		      (void *) &cur_node.tgt_offset, sizeof (void *), cbuf);
 }
@@ -638,7 +648,8 @@  attribute_hidden void
 gomp_attach_pointer (struct gomp_device_descr *devicep,
 		     struct goacc_asyncqueue *aq, splay_tree mem_map,
 		     splay_tree_key n, uintptr_t attach_to, size_t bias,
-		     struct gomp_coalesce_buf *cbufp)
+		     struct gomp_coalesce_buf *cbufp,
+		     bool allow_zero_length_array_sections)
 {
   struct splay_tree_key_s s;
   size_t size, idx;
@@ -690,11 +701,21 @@  gomp_attach_pointer (struct gomp_device_descr *devicep,
 
       if (!tn)
 	{
-	  gomp_mutex_unlock (&devicep->lock);
-	  gomp_fatal ("pointer target not mapped for attach");
+	  if (allow_zero_length_array_sections)
+	    {
+	      /* When allowing attachment to zero-length array sections, we
+		 allow attaching to NULL pointers when the target region is not
+		 mapped.  */
+	      data = 0;
+	    }
+	  else
+	    {
+	      gomp_mutex_unlock (&devicep->lock);
+	      gomp_fatal ("pointer target not mapped for attach");
+	    }
 	}
-
-      data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+      else
+	data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
 
       gomp_debug (1,
 		  "%s: attaching host %p, target %p (struct base %p) to %p\n",
@@ -950,7 +971,9 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	  has_firstprivate = true;
 	  continue;
 	}
-      else if ((kind & typemask) == GOMP_MAP_ATTACH)
+      else if ((kind & typemask) == GOMP_MAP_ATTACH
+	       || ((kind & typemask)
+		   == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
 	{
 	  tgt->list[i].key = NULL;
 	  has_firstprivate = true;
@@ -1197,7 +1220,7 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 				      (uintptr_t) *(void **) hostaddrs[j],
 				      k->tgt_offset + ((uintptr_t) hostaddrs[j]
 						       - k->host_start),
-				      sizes[j], cbufp);
+				      sizes[j], cbufp, false);
 		  }
 	      }
 	    i = j - 1;
@@ -1325,6 +1348,7 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		  ++i;
 		continue;
 	      case GOMP_MAP_ATTACH:
+	      case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 		{
 		  cur_node.host_start = (uintptr_t) hostaddrs[i];
 		  cur_node.host_end = cur_node.host_start + sizeof (void *);
@@ -1341,9 +1365,12 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 			 structured/dynamic reference counts ('n->refcount',
 			 'n->dynamic_refcount').  */
 
+		      bool zlas
+			= ((kind & typemask)
+			   == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
 		      gomp_attach_pointer (devicep, aq, mem_map, n,
 					   (uintptr_t) hostaddrs[i], sizes[i],
-					   cbufp);
+					   cbufp, zlas);
 		    }
 		  else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
 		    {
@@ -1453,9 +1480,12 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 					k->host_end - k->host_start, cbufp);
 		    break;
 		  case GOMP_MAP_POINTER:
-		    gomp_map_pointer (tgt, aq,
-				      (uintptr_t) *(void **) k->host_start,
-				      k->tgt_offset, sizes[i], cbufp);
+		  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+		    gomp_map_pointer
+		      (tgt, aq, (uintptr_t) *(void **) k->host_start,
+		       k->tgt_offset, sizes[i], cbufp,
+		       ((kind & typemask)
+			== GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION));
 		    break;
 		  case GOMP_MAP_TO_PSET:
 		    gomp_copy_host2dev (devicep, aq,
@@ -1496,7 +1526,7 @@  gomp_map_vars_internal (struct gomp_device_descr *devicep,
 					      k->tgt_offset
 					      + ((uintptr_t) hostaddrs[j]
 						 - k->host_start),
-					      sizes[j], cbufp);
+					      sizes[j], cbufp, false);
 			  }
 			}
 		    i = j - 1;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
new file mode 100644
index 00000000000..d4f9ff3e983
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -0,0 +1,34 @@ 
+extern "C" void abort ();
+
+struct S
+{
+  int *data;
+};
+
+int
+main (void)
+{
+  #define SZ 10
+  S *s = new S ();
+  s->data = new int[SZ];
+
+  for (int i = 0; i < SZ; i++)
+    s->data[i] = 0;
+
+  #pragma omp target enter data map(to: s)
+  #pragma omp target enter data map(to: s->data[:SZ])
+  #pragma omp target
+  {
+    for (int i = 0; i < SZ; i++)
+      s->data[i] = i;
+  }
+  #pragma omp target exit data map(from: s->data[:SZ])
+  #pragma omp target exit data map(from: s)
+
+  for (int i = 0; i < SZ; i++)
+    if (s->data[i] != i)
+      abort ();
+
+  return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
new file mode 100644
index 00000000000..06c6470b4ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
@@ -0,0 +1,86 @@ 
+#include <cstdlib>
+#include <cstring>
+
+template <typename L>
+void
+omp_target_loop (int begin, int end, L loop)
+{
+  #pragma omp target teams distribute parallel for
+  for (int i = begin; i < end; i++)
+    loop (i);
+}
+
+struct S
+{
+  int a, len;
+  int *ptr;
+
+  auto merge_data_func (int *iptr, int &b)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  mapped = (ptr != NULL && iptr != NULL);
+	  if (mapped)
+	    {
+	      for (int i = 0; i < len; i++)
+		ptr[i] += a + b + iptr[i];
+	    }
+	}
+	return mapped;
+      };
+    return fn;
+  }
+};
+
+int x = 1;
+
+int main (void)
+{
+  const int N = 10;
+  int *data1 = new int[N];
+  int *data2 = new int[N];
+  memset (data1, 0xab, sizeof (int) * N);
+  memset (data1, 0xcd, sizeof (int) * N);
+
+  int val = 1;
+  int &valref = val;
+  #pragma omp target enter data map(alloc: data1[:N], data2[:N])
+
+  omp_target_loop (0, N, [=](int i) { data1[i] = val; });
+  omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
+
+  #pragma omp target update from(data1[:N], data2[:N])
+
+  for (int i = 0; i < N; i++)
+    {
+      if (data1[i] != 1) abort ();
+      if (data2[i] != 2) abort ();
+    }
+
+  #pragma omp target exit data map(delete: data1[:N], data2[:N])
+
+  int b = 8;
+  S s = { 4, N, data1 };
+  auto f = s.merge_data_func (data2, b);
+
+  if (f ()) abort ();
+
+  #pragma omp target enter data map(to: data1[:N])
+  if (f ()) abort ();
+
+  #pragma omp target enter data map(to: data2[:N])
+  if (!f ()) abort ();
+
+  #pragma omp target exit data map(from: data1[:N], data2[:N])
+
+  for (int i = 0; i < N; i++)
+    {
+      if (data1[i] != 0xf) abort ();
+      if (data2[i] != 2) abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
new file mode 100644
index 00000000000..1d3561ffbd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
@@ -0,0 +1,30 @@ 
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C
new file mode 100644
index 00000000000..a591ea4c564
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-1.C
@@ -0,0 +1,29 @@ 
+extern "C" void abort ();
+
+struct S
+{
+  int a, b, c, d;
+
+  int sum (void)
+  {
+    int val = 0;
+    val += a + b + this->c + this->d;
+    return val;
+  }
+
+  int sum_offload (void)
+  {
+    int val = 0;
+    #pragma omp target map(val)
+    val += a + b + this->c + this->d;
+    return val;
+  }
+};
+
+int main (void)
+{
+  S s = { 1, 2, 3, 4 };
+  if (s.sum () != s.sum_offload ())
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C
new file mode 100644
index 00000000000..8119be8c2c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-2.C
@@ -0,0 +1,47 @@ 
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+
+extern "C" void abort ();
+
+struct T
+{
+  int x, y;
+
+  auto sum_func (int n)
+  {
+    auto fn = [=](int m) -> int
+      {
+	int v;
+	v = (x + y) * n + m;
+	return v;
+      };
+    return fn;
+  }
+
+  auto sum_func_offload (int n)
+  {
+    auto fn = [=](int m) -> int
+      {
+	int v;
+	#pragma omp target map(from:v)
+	v = (x + y) * n + m;
+	return v;
+      };
+    return fn;
+  }
+
+};
+
+int main (void)
+{
+  T a = { 1, 2 };
+
+  auto s1 = a.sum_func (3);
+  auto s2 = a.sum_func_offload (3);
+
+  if (s1 (1) != s2 (1))
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C
new file mode 100644
index 00000000000..e15f69a1623
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-3.C
@@ -0,0 +1,99 @@ 
+#include <stdio.h>
+#include <string.h>
+extern "C" void abort ();
+
+struct S
+{
+  int * ptr;
+  int ptr_len;
+
+  int *&refptr;
+  int refptr_len;
+
+  bool set_ptr (int n)
+  {
+    bool mapped;
+    #pragma omp target map(from:mapped)
+    {
+      if (ptr != NULL)
+	for (int i = 0; i < ptr_len; i++)
+	  ptr[i] = n;
+      mapped = (ptr != NULL);
+    }
+    return mapped;
+  }
+
+  bool set_refptr (int n)
+  {
+    bool mapped;
+    #pragma omp target map(from:mapped)
+    {
+      if (refptr != NULL)
+	for (int i = 0; i < refptr_len; i++)
+	  refptr[i] = n;
+      mapped = (refptr != NULL);
+    }
+    return mapped;
+  }
+};
+
+int main (void)
+{
+  #define N 10
+  int *ptr1 = new int[N];
+  int *ptr2 = new int[N];
+
+  memset (ptr1, 0, sizeof (int) * N);
+  memset (ptr2, 0, sizeof (int) * N);
+
+  S s = { ptr1, N, ptr2, N };
+
+  bool mapped;
+  int val = 123;
+
+  mapped = s.set_ptr (val);
+  if (mapped)
+    abort ();
+  if (s.ptr != ptr1)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 0)
+      abort ();
+
+  mapped = s.set_refptr (val);
+  if (mapped)
+    abort ();
+  if (s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 0)
+      abort ();
+
+  #pragma omp target data map(ptr1[:N])
+  mapped = s.set_ptr (val);
+
+  if (!mapped)
+    abort ();
+  if (s.set_refptr (0))
+    abort ();
+  if (s.ptr != ptr1 || s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != val)
+      abort ();
+
+  #pragma omp target data map(ptr2[:N])
+  mapped = s.set_refptr (val);
+
+  if (!mapped)
+    abort ();
+  if (s.set_ptr (0))
+    abort ();
+  if (s.ptr != ptr1 || s.refptr != ptr2)
+    abort ();
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != val)
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C
new file mode 100644
index 00000000000..9f53677a240
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-4.C
@@ -0,0 +1,104 @@ 
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+#include <cstdlib>
+#include <cstring>
+
+struct T
+{
+  int *ptr;
+  int ptr_len;
+
+  int *&refptr;
+  int refptr_len;
+
+  auto set_ptr_func (int n)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  if (ptr)
+	    for (int i = 0; i < ptr_len; i++)
+	      ptr[i] = n;
+	  mapped = (ptr != NULL);
+	}
+	return mapped;
+      };
+    return fn;
+  }
+
+  auto set_refptr_func (int n)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped;
+	#pragma omp target map(from:mapped)
+	{
+	  if (refptr)
+	    for (int i = 0; i < refptr_len; i++)
+	      refptr[i] = n;
+	  mapped = (refptr != NULL);
+	}
+	return mapped;
+      };
+    return fn;
+  }
+};
+
+int main (void)
+{
+  #define N 10
+  int *ptr1 = new int[N];
+  int *ptr2 = new int[N];
+
+  memset (ptr1, 0, sizeof (int) * N);
+  memset (ptr2, 0, sizeof (int) * N);
+
+  T a = { ptr1, N, ptr2, N };
+
+  auto p1 = a.set_ptr_func (1);
+  auto r2 = a.set_refptr_func (2);
+
+  if (p1 ())
+    abort ();
+  if (r2 ())
+    abort ();
+
+  if (a.ptr != ptr1)
+    abort ();
+  if (a.refptr != ptr2)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 0)
+      abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 0)
+      abort ();
+
+  #pragma omp target data map(ptr1[:N], ptr2[:N])
+  {
+    if (!p1 ())
+      abort ();
+    if (!r2 ())
+      abort ();
+  }
+
+  if (a.ptr != ptr1)
+    abort ();
+  if (a.refptr != ptr2)
+    abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr1[i] != 1)
+      abort ();
+
+  for (int i = 0; i < N; i++)
+    if (ptr2[i] != 2)
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C
new file mode 100644
index 00000000000..e71c566687d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-5.C
@@ -0,0 +1,30 @@ 
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+  T a, b, c, d;
+
+  T sum (void)
+  {
+    T val = 0;
+    val += a + b + this->c + this->d;
+    return val;
+  }
+
+  T sum_offload (void)
+  {
+    T val = 0;
+    #pragma omp target map(val)
+    val += a + b + this->c + this->d;
+    return val;
+  }
+};
+
+int main (void)
+{
+  S<int> s = { 1, 2, 3, 4 };
+  if (s.sum () != s.sum_offload ())
+    abort ();
+  return 0;
+}