[PATCH, OpenMP 5.0] Improve OpenMP target support for C++ (includes PR92120 v3)

Chung-Lin Tang cltang@codesourcery.com
Thu May 20 15:04:18 GMT 2021


Hi Jakub,
the attached patch is a combination of the below patches already pushed to devel/omp/gcc-10,
some are kind of transient bug fixes, but listing all for completeness:

aadfc984: [PATCH] Target mapping C++ members inside member functions
https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562467.html

36a1ebdb: [PATCH] OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120)
https://gcc.gnu.org/pipermail/gcc-patches/2020-November/558975.html

bf8605f1: [PATCH] Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF (INDIRECT_REF ...)) map clauses.
https://gcc.gnu.org/pipermail/gcc-patches/2021-February/564976.html

da047f63: [PATCH] Fix regression of array members in OpenMP map clauses.
https://gcc.gnu.org/pipermail/gcc-patches/2021-March/566086.html

4e714eaa: [PATCH] Fix template case of non-static member access inside member functions
https://gcc.gnu.org/pipermail/gcc-patches/2021-March/566592.html

2ed80263: [PATCH] Lambda capturing of pointers and references in target directives
https://gcc.gnu.org/pipermail/gcc-patches/2021-March/566935.html

08caada8: Arrow operator handling for C front-end in OpenMP map clauses
https://gcc.gnu.org/pipermail/gcc-patches/2021-March/566419.html

To summarize, this patch set is an improvement for OpenMP target support for C++,
including for inside non-static members, lambda objects, and struct member deref access expressions.
The corresponding modifications for the C front-end are also included.

This patch supercedes the prior versions of my PR92120 patch (implicit C++ map(this[:1])),
so dubbing this "v3" of patch for that PR.

Prior versions of the PR92120 patch was implemented by recording uses of 'this' in the parser,
and then use the recorded uses during "finish" to create the implicit maps.

When working on supporting lambda objects, this required using a tree-walk style processing of
the OMP_TARGET body, so in only made sense to merge the entire 'this' processing together with it,
so a large part of the parser changes were dropped, with the main processing in semantics.c now.

Other parser changes to support '->' in map clauses are also with this patch.

Tested without regressions on x86_64-linux with nvptx offloading, okay for trunk?

Thanks,
Chung-Lin

2021-05-20  Chung-Lin Tang  <cltang@codesourcery.com>

	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-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-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 b9930d487fd..5f6d3184e1a 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15749,7 +15749,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 c3b85da1125..8f32d808c63 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13092,6 +13092,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))
@@ -13928,6 +13934,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;
@@ -14659,6 +14666,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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14715,6 +14727,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_)
 	    {
@@ -14787,6 +14808,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
@@ -14836,8 +14858,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
 	  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 aa202715873..2785d9d3659 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7591,6 +7591,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 bc0505df502..f25833f0761 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38011,7 +38011,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);
@@ -42030,8 +42030,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);
@@ -42135,15 +42133,9 @@ 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];
-	  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;
+	  finish_omp_target (pragma_tok->location,
+			     cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
+	  return true;
 	}
       else if (!flag_openmp)  /* flag_openmp_simd  */
 	{
@@ -42180,49 +42172,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 cbd2f3dc338..1d9cf4a97ff 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18898,6 +18898,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 0d590c318fb..f45719bc7b1 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4935,6 +4935,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);
@@ -4963,8 +4966,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);
@@ -4984,6 +4991,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)
@@ -5498,6 +5508,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))
@@ -6511,6 +6523,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
@@ -7605,6 +7618,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 (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -7667,10 +7685,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)
@@ -7719,6 +7742,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)
@@ -7745,7 +7776,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%>"
@@ -7758,6 +7791,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
@@ -8865,6 +8899,475 @@ 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;
+};
+
+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;
+}
+
+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;
+	  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;
+  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 719a4e16391..ee1555f2e32 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"
@@ -8543,7 +8544,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;
@@ -9012,7 +9013,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)
@@ -9028,12 +9036,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
@@ -9044,6 +9055,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))
@@ -9052,6 +9064,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))
@@ -9063,8 +9076,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
@@ -9134,7 +9150,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
@@ -9191,7 +9210,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)
@@ -9217,7 +9239,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);
@@ -9228,7 +9254,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)
@@ -9236,7 +9273,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)
 			{
@@ -9270,7 +9308,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)
 		    {
@@ -9379,6 +9442,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;
@@ -9406,6 +9476,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
@@ -10659,6 +10747,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
@@ -10705,10 +10799,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 d1136d181b3..6f98e589810 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12019,6 +12019,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-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 0a575eb9dad..eaf273b2616 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -816,6 +816,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:
@@ -894,6 +895,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 ();
 	}
@@ -912,6 +916,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;
@@ -919,6 +926,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 ef1bb4907b6..931dd5b1feb 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1196,7 +1196,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 405574dfa2b..7c0ab7d9d6e 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -939,7 +939,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);
 }
@@ -1143,7 +1143,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').  */
 	    }
@@ -1161,7 +1161,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 2150e5d79b2..f152e1c11b8 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -371,7 +371,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
 			unsigned char kind, bool always_to_flag,
 			struct gomp_coalesce_buf *cbuf)
 {
-  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);
@@ -412,7 +413,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;
@@ -434,16 +436,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);
 }
@@ -514,7 +524,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;
@@ -566,11 +577,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",
@@ -824,7 +845,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;
@@ -1070,7 +1093,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;
@@ -1197,6 +1220,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 *);
@@ -1213,9 +1237,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)
 		    {
@@ -1298,9 +1325,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,
@@ -1335,7 +1365,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-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;
+}


More information about the Gcc-patches mailing list