Lambda capturing of pointers and references in target directives

This patch implements proper lambda capturing of pointer and reference variables
as specified in OpenMP 5.0. We map the entire closure object as a to-map,
attach pointers to zero-length array sections, and perform mapping of
references.

2021-03-18  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/cp/ChangeLog:

	* cp-tree.h (set_omp_target_this_expr): Delete.
	(finish_omp_target_clauses): New prototype.
	* lambda.c (lambda_expr_this_capture): Remove call to
	set_omp_target_this_expr.
	* parser.c (cp_parser_omp_target): Likewise.
	* pt.c (tsubst_expr): Add call to finish_omp_target_clauses for target
	directives.
	* semantics.c (omp_target_this_expr): Delete.
	(omp_target_ptr_members_accessed): Delete.
	(finish_non_static_data_member): Remove call to
	set_omp_target_this_expr. Remove use of omp_target_ptr_members_accessed.
	(finish_this_expr): Remove call to set_omp_target_this_expr.
	(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, with code factored out from
	finish_omp_target. Add lambda object handling case.
	(finish_omp_target): Factor code out and adjust to use
	finish_omp_target_clauses.
	(finish_omp_clauses): Revert prior "Adjustments to allow '*ptr' and
	'ptr->member' cases in map clausess.", since not needed with new
	organization of target-directive clause processing.

gcc/testsuite/ChangeLog:

	* g++.dg/gomp/target-lambda-1.C: New test.

libgomp/testsuite/ChangeLog:

	* libgomp.c++/target-lambda-1.C: New test.

7 files changed, 518 insertions, 168 deletions

diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index b77bdc380a0..247a3bb1ec3 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7316,7 +7316,7 @@ extern void finish_lambda_scope			(void);
 extern tree start_lambda_function		(tree fn, tree lambda_expr);
 extern void finish_lambda_function		(tree body);
 extern tree finish_omp_target			(location_t, tree, tree, bool);
-extern void set_omp_target_this_expr		(tree);
+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/lambda.c b/gcc/cp/lambda.c
index 9ecf0dbed0c..b55c2f85d27 100644
--- a/gcc/cp/lambda.c
+++ b/gcc/cp/lambda.c
@@ -842,9 +842,6 @@ lambda_expr_this_capture (tree lambda, int add_capture_p)
 	 type cast (_expr.cast_ 5.4) to the type of 'this'. [ The cast
 	 ensures that the transformed expression is an rvalue. ] */
       result = rvalue (result);
-
-      /* Acknowledge to OpenMP target that 'this' was referenced.  */
-      set_omp_target_this_expr (result);
     }
 
   return result;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 1af233690a2..9fc2a9b05eb 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40786,7 +40786,6 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 	  keep_next_level (true);
 	  tree sb = begin_omp_structured_block (), ret;
 	  unsigned save = cp_parser_begin_omp_structured_block (parser);
-	  set_omp_target_this_expr (NULL_TREE);
 	  switch (ccode)
 	    {
 	    case OMP_TEAMS:
@@ -40881,7 +40880,6 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 					    "#pragma omp target", pragma_tok);
   c_omp_adjust_map_clauses (clauses, true);
   keep_next_level (true);
-  set_omp_target_this_expr (NULL_TREE);
   tree body = cp_parser_omp_structured_block (parser, if_p);
 
   finish_omp_target (pragma_tok->location, clauses, body, false);
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 90cee31bb5a..139d1075986 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18631,6 +18631,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 5b62fa35eb8..3e290767d5c 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -61,11 +61,6 @@ static hash_map<tree, tree> *omp_private_member_map;
 static vec<tree> omp_private_member_vec;
 static bool omp_private_member_ignore_next;
 
-/* Used for OpenMP target region 'this' references.  */
-static tree omp_target_this_expr = NULL_TREE;
-
-static hash_map<tree, tree> omp_target_ptr_members_accessed;
-
 /* Deferred Access Checking Overview
    ---------------------------------
 
@@ -1896,7 +1891,6 @@ tree
 finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
 {
   gcc_assert (TREE_CODE (decl) == FIELD_DECL);
-  tree orig_object = object;
   bool try_omp_private = !object && omp_private_member_map;
   tree ret;
 
@@ -1935,14 +1929,6 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
       return error_mark_node;
     }
 
-  if (orig_object == NULL_TREE)
-    {
-      tree this_expr = TREE_OPERAND (object, 0);
-
-      /* Acknowledge to OpenMP target that 'this' was referenced.  */
-      set_omp_target_this_expr (this_expr);
-    }
-
   if (current_class_ptr)
     TREE_USED (current_class_ptr) = 1;
   if (processing_template_decl)
@@ -2003,13 +1989,6 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
       if (v)
 	ret = convert_from_reference (*v);
     }
-  else if (omp_target_this_expr
-	   && TREE_TYPE (ret)
-	   && POINTER_TYPE_P (TREE_TYPE (ret)))
-    {
-      if (omp_target_ptr_members_accessed.get (decl) == NULL)
-	omp_target_ptr_members_accessed.put (decl, ret);
-    }
 
   return ret;
 }
@@ -2773,9 +2752,6 @@ finish_this_expr (void)
       /* The keyword 'this' is a prvalue expression.  */
       result = rvalue (result);
 
-      /* Acknowledge to OpenMP target that 'this' was referenced.  */
-      set_omp_target_this_expr (result);
-
       return result;
     }
 
@@ -6407,7 +6383,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bool order_seen = false;
   bool schedule_seen = false;
   bool oacc_async = false;
-  bool indirect_ref_p = false;
   bool indir_component_ref_p = false;
   tree last_iterators = NULL_TREE;
   bool last_iterators_remove = false;
@@ -7517,14 +7492,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      indir_component_ref_p = true;
 	      STRIP_NOPS (t);
 	    }
-	  indirect_ref_p = false;
-	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
-	      && INDIRECT_REF_P (t))
-	    {
-	      t = TREE_OPERAND (t, 0);
-	      indirect_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)
@@ -7560,12 +7527,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      break;
 		    }
 		  t = TREE_OPERAND (t, 0);
-		  if (INDIRECT_REF_P (t))
-		    {
-		      t = TREE_OPERAND (t, 0);
-		      indir_component_ref_p = true;
-		      STRIP_NOPS (t);
-		    }
 		}
 	      if (remove)
 		break;
@@ -7629,7 +7590,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		       || (OMP_CLAUSE_MAP_KIND (c)
 			   != GOMP_MAP_FIRSTPRIVATE_POINTER))
 		   && !indir_component_ref_p
-		   && !indirect_ref_p
 		   && !cxx_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -7714,8 +7674,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  else
 	    {
-	      if (!indirect_ref_p && !indir_component_ref_p)
-		bitmap_set_bit (&map_head, DECL_UID (t));
+	      bitmap_set_bit (&map_head, DECL_UID (t));
 	      if (t != OMP_CLAUSE_DECL (c)
 		  && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
 		bitmap_set_bit (&map_field_head, DECL_UID (t));
@@ -8683,26 +8642,126 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
   return add_stmt (stmt);
 }
 
-void
-set_omp_target_this_expr (tree this_val)
+/* Used to walk OpenMP target directive body.  */
+
+struct omp_target_walk_data
 {
-  omp_target_this_expr = this_val;
+  tree current_object;
+  bool this_expr_accessed;
+
+  hash_map<tree, tree> ptr_members_accessed;
+  hash_set<tree> lambda_objects_accessed;
 
-  if (omp_target_this_expr == NULL_TREE)
-    omp_target_ptr_members_accessed.empty ();
+  tree current_closure;
+  hash_set<tree> closure_vars_accessed;
+};
+
+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_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))
+	data->lambda_objects_accessed.add (t);
+      *walk_subtrees = 0;
+      return NULL_TREE;
+    }
+
+  return NULL_TREE;
 }
 
-tree
-finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
+void
+finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
 {
-  tree last_inserted_clause = NULL_TREE;
+  omp_target_walk_data data;
+  data.this_expr_accessed = false;
 
-  if (omp_target_this_expr)
+  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; *c; c = &OMP_CLAUSE_CHAIN (*c))
+      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),
@@ -8722,23 +8781,72 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	  OMP_CLAUSE_DECL (c)
 	    = build_indirect_ref (loc, closure, RO_UNARY_STAR);
 	  OMP_CLAUSE_SIZE (c)
-	    = (processing_template_decl
-	       ? NULL_TREE
-	       : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))));
+	    = 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;
-	  OMP_CLAUSE_CHAIN (c2) = clauses;
-	  OMP_CLAUSE_CHAIN (c) = c2;
-	  last_inserted_clause = c2;
-	  clauses = c;
+	  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.  */
@@ -8753,12 +8861,13 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	      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.  */
-	      *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc);
-	      OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c2);
-	      OMP_CLAUSE_CHAIN (c2) = this_map;
-	      last_inserted_clause = nc;
+	      new_clauses.safe_push (this_map);
+	      new_clauses.safe_push (nc);
 	    }
 	  else
 	    {
@@ -8767,9 +8876,7 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	      OMP_CLAUSE_DECL (c3)
 		= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
 	      OMP_CLAUSE_SIZE (c3)
-		= (processing_template_decl
-		   ? NULL_TREE
-		   : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
+		= 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);
@@ -8777,10 +8884,8 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	      OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
 	      OMP_CLAUSE_SIZE (c4) = size_zero_node;
 
-	      OMP_CLAUSE_CHAIN (c3) = c4;
-	      OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2);
-	      OMP_CLAUSE_CHAIN (c2) = c3;
-	      last_inserted_clause = c4;
+	      new_clauses.safe_push (c3);
+	      new_clauses.safe_push (c4);
 	    }
 	}
       else
@@ -8794,112 +8899,177 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	      OMP_CLAUSE_DECL (c)
 		= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
 	      OMP_CLAUSE_SIZE (c)
-		= (processing_template_decl
-		   ? NULL_TREE
-		   : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
+		= 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;
-	      OMP_CLAUSE_CHAIN (c2) = clauses;
-	      OMP_CLAUSE_CHAIN (c) = c2;
-	      clauses = c;
-	      last_inserted_clause = c2;
+
+	      new_clauses.safe_push (c);
+	      new_clauses.safe_push (c2);
 	    }
 	}
-      omp_target_this_expr = NULL_TREE;
-    }
-
-  if (last_inserted_clause && !omp_target_ptr_members_accessed.is_empty ())
-    for (hash_map<tree, tree>::iterator i
-	   = omp_target_ptr_members_accessed.begin ();
-	 i != omp_target_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 nc = OMP_CLAUSE_CHAIN (last_inserted_clause);
-	     nc != NULL_TREE; nc = OMP_CLAUSE_CHAIN (nc))
+      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)
 	  {
-	    /* If map(this->ptr[:N] already exists, avoid creating another
-	       such map.  */
-	    tree decl = OMP_CLAUSE_DECL (nc);
-	    if ((TREE_CODE (decl) == INDIRECT_REF
-		 || TREE_CODE (decl) == MEM_REF)
-		&& operand_equal_p (TREE_OPERAND (decl, 0),
-				    ptr_member))
-	      goto next_ptr_member;
-	  }
+	    /* 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;
 
-	if (!cxx_mark_addressable (ptr_member))
-	  gcc_unreachable ();
+	    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 (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;
-
-	    OMP_CLAUSE_CHAIN (c) = c2;
-	    OMP_CLAUSE_CHAIN (c2) = c3;
-	    OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last_inserted_clause);
-
-	    OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
-	    last_inserted_clause = 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;
-
-	    OMP_CLAUSE_CHAIN (c) = c2;
-	    OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last_inserted_clause);
-	    OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
-	    last_inserted_clause = c2;
+	    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:
+	    ;
 	  }
-	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;
+	  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_TO);
+		  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;
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/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;
+}