openmp: Improve OpenMP target support for C++ (PR92120)

This patch implements several C++ specific mapping capabilities introduced for
OpenMP 5.0, including implicit mapping of this[:1] for non-static member
functions, zero-length array section mapping of pointer-typed members,
lambda captured variable access in target regions, and use of lambda objects
inside target regions.

Several adjustments to the C/C++ front-ends to allow more member-access syntax
as valid is also included.

	PR middle-end/92120

gcc/cp/ChangeLog:

	* 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. Remove case of rejecting
	'this' when not in declare simd.
	(handle_omp_array_sections): Likewise.
	(finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP
	map clauses. Handle 'A->member' case in map clauses. Remove case of
	rejecting 'this' when not in declare simd.
	(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/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* 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.

11 files changed, 516 insertions, 28 deletions

diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 299cf42..832ca6e 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1283,7 +1283,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 5988db0..82d8dac 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 5d3103a..6b02daf 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -543,7 +543,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);
@@ -616,7 +617,8 @@ get_implicit (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;
@@ -638,16 +640,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 *),
 		      true, cbuf);
@@ -724,7 +734,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;
@@ -776,11 +787,19 @@ 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",
@@ -1038,7 +1057,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;
@@ -1287,7 +1308,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;
@@ -1416,6 +1437,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 *);
@@ -1432,9 +1454,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)
 		    {
@@ -1545,9 +1570,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 					false, 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,
@@ -1589,7 +1617,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 0000000..d4f9ff3
--- /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 0000000..06c6470
--- /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 0000000..1d3561f
--- /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 0000000..a591ea4
--- /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 0000000..8119be8
--- /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 0000000..e15f69a
--- /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 0000000..9f53677
--- /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 0000000..e71c566
--- /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;
+}