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

Chung-Lin Tang cltang@codesourcery.com
Fri Jun 18 14:25:16 GMT 2021


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.

-------------- next part --------------
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;
+}
-------------- next part --------------
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 ();
+}


More information about the Gcc-patches mailing list