OpenMP/C++: Fix (first)private clause with member variables [PR110347]

OpenMP permits '(first)private' for C++ member variables, which GCC handles
by tagging those by DECL_OMP_PRIVATIZED_MEMBER, adding a temporary VAR_DECL
and DECL_VALUE_EXPR pointing to the 'this->member_var' in the C++ front end.

The idea is that in omp-low.cc, the DECL_VALUE_EXPR is used before the
region (for 'firstprivate'; ignored for 'private') while in the region,
the DECL itself is used.

In gimplify, the value expansion is suppressed and deferred if the
  lang_hooks.decls.omp_disregard_value_expr (decl, shared)
returns true - which is never the case if 'shared' is true. In OpenMP 4.5,
only 'map' and 'use_device_ptr' was permitted for the 'target' directive.
And when OpenMP 5.0's 'private'/'firstprivate' clauses was added, the
update that 'shared' is only true for 'map' was missed.

However, just enabling it for all '!shared' will cause issues with
Lambda closures ("__closure->this->...") for which also a DECL_VALUE_EXPR
exists but that is not related to DECL_OMP_PRIVATIZED_MEMBER. Solution:
Update the lang hook to take a Boolean argument, indicating whether it
is called for a target region or not.

2024-02-16  Tobias Burnus  <tburnus@baylibre.com>
	    Jakub Jelinek  <jakub@redhat.com>

	PR c++/110347

gcc/cp/ChangeLog:

	* cp-gimplify.cc (cxx_omp_disregard_value_expr): Add new
	Boolean argument and use it.
	* cp-tree.h (cxx_omp_disregard_value_expr): Update prototype.

gcc/fortran/ChangeLog:

	* trans-openmp.cc (gfc_omp_disregard_value_expr): Add
	unused Boolean argument.
	* trans.h (gfc_omp_disregard_value_expr): Update
	prototype.

gcc/ChangeLog:

	* gimplify.cc (omp_notice_variable): Update call to
	lang_hooks.decls.omp_disregard_value_expr.
	(omp_notice_variable): Likewise; fix 'shared' arg for
	(first)private in target regions.
	* hooks.cc (hook_bool_tree_bool_bool_false): New.
	* hooks.h (hook_bool_tree_bool_bool_false): New.
	* langhooks-def.h (LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR):
	Use it.
	* langhooks.h (struct lang_hooks_for_decls): Add second
	Boolean argument.
	* omp-low.cc (omp_member_access_dummy_var): Update
	lang_hooks.decls.omp_disregard_value_expr call.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-lambda-3.C: Moved from
	gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling.
	* testsuite/libgomp.c++/firstprivate-c++-1.C: New test.
	* testsuite/libgomp.c++/firstprivate-c++-2.C: New test.
	* testsuite/libgomp.c++/private-c++-1.C: New test.
	* testsuite/libgomp.c++/private-c++-2.C: New test.
	* testsuite/libgomp.c++/use_device_ptr-c++-1.C: New test.

gcc/testsuite/ChangeLog:

	* g++.dg/gomp/target-lambda-1.C: Moved to become a
	run-time test under testsuite/libgomp.c++.

Co-authored-by: Jakub Jelinek <jakub@redhat.com>

 gcc/cp/cp-gimplify.cc                              |   7 +-
 gcc/cp/cp-tree.h                                   |   2 +-
 gcc/fortran/trans-openmp.cc                        |   2 +-
 gcc/fortran/trans.h                                |   2 +-
 gcc/gimplify.cc                                    |  12 +-
 gcc/hooks.cc                                       |   6 +
 gcc/hooks.h                                        |   1 +
 gcc/langhooks-def.h                                |   2 +-
 gcc/langhooks.h                                    |   5 +-
 gcc/omp-low.cc                                     |   2 +-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C        |  94 -------
 libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C | 305 +++++++++++++++++++++
 libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C | 125 +++++++++
 libgomp/testsuite/libgomp.c++/private-c++-1.C      | 247 +++++++++++++++++
 libgomp/testsuite/libgomp.c++/private-c++-2.C      | 117 ++++++++
 libgomp/testsuite/libgomp.c++/target-lambda-3.C    | 104 +++++++
 .../testsuite/libgomp.c++/use_device_ptr-c++-1.C   | 125 +++++++++
 17 files changed, 1048 insertions(+), 110 deletions(-)

diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc
index 30e94797f9f..dcc46d86619 100644
--- a/gcc/cp/cp-gimplify.cc
+++ b/gcc/cp/cp-gimplify.cc
@@ -2754,10 +2754,11 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */)
 /* Return true if DECL's DECL_VALUE_EXPR (if any) should be
    disregarded in OpenMP construct, because it is going to be
    remapped during OpenMP lowering.  SHARED is true if DECL
-   is going to be shared, false if it is going to be privatized.  */
+   is going to be shared, false if it is going to be privatized. TARGET is
+   true if this for an OpenMP target/OpenACC compute region.   */
 
 bool
-cxx_omp_disregard_value_expr (tree decl, bool shared)
+cxx_omp_disregard_value_expr (tree decl, bool shared, bool target)
 {
   if (shared)
     return false;
@@ -2767,7 +2768,7 @@ cxx_omp_disregard_value_expr (tree decl, bool shared)
       && DECL_LANG_SPECIFIC (decl)
       && DECL_OMP_PRIVATIZED_MEMBER (decl))
     return true;
-  if (VAR_P (decl) && DECL_CONTEXT (decl) && is_capture_proxy (decl))
+  if (!target && VAR_P (decl) && DECL_CONTEXT (decl) && is_capture_proxy (decl))
     return true;
   return false;
 }
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 334c11396c2..2dc200cd43e 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -8464,7 +8464,7 @@ extern tree cxx_omp_clause_assign_op		(tree, tree, tree);
 extern tree cxx_omp_clause_dtor			(tree, tree);
 extern void cxx_omp_finish_clause		(tree, gimple_seq *, bool);
 extern bool cxx_omp_privatize_by_reference	(const_tree);
-extern bool cxx_omp_disregard_value_expr	(tree, bool);
+extern bool cxx_omp_disregard_value_expr	(tree, bool, bool);
 extern void cp_fold_function			(tree);
 extern tree cp_fold_maybe_rvalue		(tree, bool);
 extern tree cp_fold_rvalue			(tree);
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a2bf15665b3..74e213ab09e 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1791,7 +1791,7 @@ gfc_omp_scalar_target_p (tree decl)
    is going to be shared, false if it is going to be privatized.  */
 
 bool
-gfc_omp_disregard_value_expr (tree decl, bool shared)
+gfc_omp_disregard_value_expr (tree decl, bool shared, bool /* target */)
 {
   if (GFC_DECL_COMMON_OR_EQUIV (decl)
       && DECL_HAS_VALUE_EXPR_P (decl))
diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h
index 2e10ce1a9b3..d8e640ade27 100644
--- a/gcc/fortran/trans.h
+++ b/gcc/fortran/trans.h
@@ -840,7 +840,7 @@ void gfc_omp_finish_clause (tree, gimple_seq *, bool);
 bool gfc_omp_allocatable_p (tree);
 bool gfc_omp_scalar_p (tree, bool);
 bool gfc_omp_scalar_target_p (tree);
-bool gfc_omp_disregard_value_expr (tree, bool);
+bool gfc_omp_disregard_value_expr (tree, bool, bool);
 bool gfc_omp_private_debug_clause (tree, bool);
 bool gfc_omp_private_outer_ref (tree);
 struct gimplify_omp_ctx;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 7f79b3cc7e6..dc524dc12b0 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7990,7 +7990,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
      the whole block.  For C++ and Fortran, it can also be true under certain
      other conditions, if DECL_HAS_VALUE_EXPR.  */
   if (RECORD_OR_UNION_TYPE_P (type))
-    is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false);
+    is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false, true);
 
   if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
       && is_global_var (decl)
@@ -8092,7 +8092,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
     }
 
   if (ctx->region_type == ORT_NONE)
-    return lang_hooks.decls.omp_disregard_value_expr (decl, false);
+    return lang_hooks.decls.omp_disregard_value_expr (decl, false, false);
 
   if (is_global_var (decl))
     {
@@ -8148,9 +8148,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	/* For OpenACC, as remarked above, defer expansion.  */
 	shared = false;
       else
-	shared = true;
+	shared = (flags & GOVD_MAP) != 0;
 
-      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
+      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, true);
       if (n == NULL)
 	{
 	  unsigned nflags = flags;
@@ -8305,7 +8305,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
       omp_add_variable (ctx, decl, flags);
 
       shared = (flags & GOVD_SHARED) != 0;
-      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
+      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, false);
       goto do_outer;
     }
 
@@ -8350,7 +8350,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
     shared = false;
   else
     shared = ((flags | n->value) & GOVD_SHARED) != 0;
-  ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
+  ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, false);
 
   /* If nothing changed, there's nothing left to do.  */
   if ((n->value & flags) == flags)
diff --git a/gcc/hooks.cc b/gcc/hooks.cc
index 28769074222..32eccfd5f08 100644
--- a/gcc/hooks.cc
+++ b/gcc/hooks.cc
@@ -343,6 +343,12 @@ hook_bool_tree_bool_false (tree, bool)
   return false;
 }
 
+bool
+hook_bool_tree_bool_bool_false (tree, bool, bool)
+{
+  return false;
+}
+
 bool
 hook_bool_rtx_insn_true (rtx_insn *)
 {
diff --git a/gcc/hooks.h b/gcc/hooks.h
index 924748420e6..02dc63dc3c1 100644
--- a/gcc/hooks.h
+++ b/gcc/hooks.h
@@ -72,6 +72,7 @@ extern bool hook_bool_rtx_mode_int_int_intp_bool_false (rtx, machine_mode,
 extern bool hook_bool_tree_tree_false (tree, tree);
 extern bool hook_bool_tree_tree_true (tree, tree);
 extern bool hook_bool_tree_bool_false (tree, bool);
+extern bool hook_bool_tree_bool_bool_false (tree, bool, bool);
 extern bool hook_bool_wint_wint_uint_bool_true (const widest_int &,
 						const widest_int &,
 						unsigned int, bool);
diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h
index f5c67b6823c..67c100a0af3 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -263,7 +263,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
 #define LANG_HOOKS_OMP_PREDETERMINED_SHARING lhd_omp_predetermined_sharing
 #define LANG_HOOKS_OMP_PREDETERMINED_MAPPING lhd_omp_predetermined_mapping
 #define LANG_HOOKS_OMP_REPORT_DECL lhd_pass_through_t
-#define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR hook_bool_tree_bool_false
+#define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR hook_bool_tree_bool_bool_false
 #define LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE hook_bool_tree_bool_false
 #define LANG_HOOKS_OMP_PRIVATE_OUTER_REF hook_bool_tree_false
 #define LANG_HOOKS_OMP_CLAUSE_DEFAULT_CTOR hook_tree_tree_tree_tree_null
diff --git a/gcc/langhooks.h b/gcc/langhooks.h
index 5a4dfb6ef62..68bd91f3c62 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -280,8 +280,9 @@ struct lang_hooks_for_decls
   /* Return true if DECL's DECL_VALUE_EXPR (if any) should be
      disregarded in OpenMP construct, because it is going to be
      remapped during OpenMP lowering.  SHARED is true if DECL
-     is going to be shared, false if it is going to be privatized.  */
-  bool (*omp_disregard_value_expr) (tree, bool);
+     is going to be shared, false if it is going to be privatized.  TARGET
+     is true when this if for an OpenMP target/OPenACC compute contruct.  */
+  bool (*omp_disregard_value_expr) (tree, bool, bool);
 
   /* Return true if DECL that is shared iff SHARED is true should
      be put into OMP_CLAUSE_PRIVATE_DEBUG.  */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 4d003f42098..7b4631029c7 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -275,7 +275,7 @@ omp_member_access_dummy_var (tree decl)
       || !DECL_ARTIFICIAL (decl)
       || !DECL_IGNORED_P (decl)
       || !DECL_HAS_VALUE_EXPR_P (decl)
-      || !lang_hooks.decls.omp_disregard_value_expr (decl, false))
+      || !lang_hooks.decls.omp_disregard_value_expr (decl, false, false))
     return NULL_TREE;
 
   tree v = DECL_VALUE_EXPR (decl);
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
deleted file mode 100644
index 5ce8ceadb19..00000000000
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ /dev/null
@@ -1,94 +0,0 @@
-// 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.* 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\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C b/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C
new file mode 100644
index 00000000000..ae5d4fbe1bf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C
@@ -0,0 +1,305 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+struct S {
+  int A, B[10], *C;
+  void f (int dev);
+  void g (int dev);
+};
+
+template<typename T>
+struct St {
+  T A, B[10], *C;
+  void ft (int dev);
+  void gt (int dev);
+};
+
+
+void
+S::f (int dev)
+{
+  A = 5;
+  C = (int *) malloc (sizeof (int) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+                     firstprivate(c_saved) device(dev)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+void
+S::g (int dev)
+{
+  A = 5;
+  C = (int *) malloc (sizeof (int) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+                      allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+                      device(dev)
+    {
+#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+#endif
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \
+                       allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+
+template<typename T>
+void
+St<T>::ft (int dev)
+{
+  A = 5;
+  C = (T *) malloc (sizeof (T) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+                     firstprivate(c_saved) device(dev)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+template<typename T>
+void
+St<T>::gt (int dev)
+{
+  A = 5;
+  C = (T *) malloc (sizeof (T) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \
+                     allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+                     device(dev)
+    {
+#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+#endif
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \
+                       allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+    {
+      if (A != 5)
+	abort ();
+      for (int i = 0; i < 10; i++)
+	if (B[i] != i + 5)
+	  abort ();
+      if (c_saved != (uintptr_t) C)
+	abort ();
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+int
+main ()
+{
+  struct S s;
+  struct St<int> st;
+  for (int dev = 0; dev <= omp_get_num_devices(); dev++)
+    {
+      s.f (dev);
+      st.ft (dev);
+      s.g (dev);
+      st.gt (dev);
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C b/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C
new file mode 100644
index 00000000000..a4f2514b591
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C
@@ -0,0 +1,125 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+
+struct t {
+  int A;
+  void f (int dev);
+};
+
+void
+t::f (int dev)
+{
+  int B = 49;
+
+  A = 7;
+  #pragma omp parallel firstprivate(A) if(0) shared(B) default(none)
+  {
+    if (A != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", A); __builtin_abort (); }
+    A = 5;
+    B = A;
+  }
+  if (A != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", A); __builtin_abort (); }
+  if (B != 5) { __builtin_printf("ERROR 1a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+  #pragma omp parallel firstprivate(A)if(0) shared(B) default(none)
+  {
+    if (A != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", A); __builtin_abort (); }
+    A = 6;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", A); __builtin_abort (); }
+  if (B != 6) { __builtin_printf("ERROR 2a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+
+  #pragma omp target firstprivate(A) map(from:B) device(dev)
+  {
+    if (A != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", A); __builtin_abort (); }
+    A = 7;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", A); __builtin_abort (); }
+  if (B != 7) { __builtin_printf("ERROR 3a: %d\n", B); __builtin_abort (); }
+  A = 9; B = 49;
+  #pragma omp target firstprivate(A) map(from:B) device(dev)
+  {
+    if (A != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", A); __builtin_abort (); }
+    A = 8;
+    B = A;
+  }
+  if (A != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", A); __builtin_abort (); }
+  if (B != 8) { __builtin_printf("ERROR 4a: %d\n", B); __builtin_abort (); }
+}
+
+
+template <typename T>
+struct tt {
+  T C;
+  void g (int dev);
+};
+
+template <typename T>
+void
+tt<T>::g (int dev)
+{
+  T D = 49;
+  C = 7;
+  #pragma omp parallel firstprivate(C) if(0) shared(D) default(none)
+  {
+    if (C != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", C);__builtin_abort (); }
+    C = 5;
+    D = C;
+  }
+  if (C != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", C);__builtin_abort (); }
+  if (D != 5) { __builtin_printf("ERROR 1a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp parallel firstprivate(C)if(0) shared(D) default(none)
+  {
+    if (C != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", C);__builtin_abort (); }
+    C = 6;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", C);__builtin_abort (); }
+  if (D != 6) { __builtin_printf("ERROR 2a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev)
+  {
+    if (C != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", C);__builtin_abort (); }
+    C = 7;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", C);__builtin_abort (); }
+  if (D != 7) { __builtin_printf("ERROR 3a: %d\n", D);__builtin_abort (); }
+  C = 9; D = 49;
+  #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev)
+  {
+    if (C != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", C);__builtin_abort (); }
+    C = 8;
+    D = C;
+  }
+  if (C != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", C); __builtin_abort (); }
+  if (D != 8) { __builtin_printf("ERROR 4a: %d\n", D); }
+}
+
+void
+foo ()
+{
+  struct t x;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    x.f (dev);
+}
+
+void
+bar ()
+{
+  struct tt<int> y;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    y.g (dev);
+}
+
+int
+main ()
+{
+  foo ();
+  bar ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/private-c++-1.C b/libgomp/testsuite/libgomp.c++/private-c++-1.C
new file mode 100644
index 00000000000..19ee726a222
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/private-c++-1.C
@@ -0,0 +1,247 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+struct S {
+  int A, B[10], *C;
+  void f (int dev);
+  void g (int dev);
+};
+
+template<typename T>
+struct St {
+  T A, B[10], *C;
+  void ft (int dev);
+  void gt (int dev);
+};
+
+
+void
+S::f (int dev)
+{
+  A = 5;
+  C = (int *) malloc (sizeof (int) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target private(A) private(B) private(C) device(dev)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) private(A) private(B) private(C)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+void
+S::g (int dev)
+{
+  A = 5;
+  C = (int *) malloc (sizeof (int) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target private(A) private(B) private(C) \
+                     allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+                     device(dev)
+    {
+#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+#endif
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) private(A) private(B) private(C) \
+                       allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+    {
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+
+template<typename T>
+void
+St<T>::ft (int dev)
+{
+  A = 5;
+  C = (T *) malloc (sizeof (T) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target private(A) private(B) private(C) device(dev)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) private(A) private(B) private(C)
+    {
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+template<typename T>
+void
+St<T>::gt (int dev)
+{
+  A = 5;
+  C = (T *) malloc (sizeof (T) * 10);
+  uintptr_t c_saved = (uintptr_t) C;
+  for (int i = 0; i < 10; i++)
+    B[i] = C[i] = i+5;
+
+  #pragma omp target private(A) private(B) private(C) \
+                     allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \
+                     device(dev)
+    {
+#if 0  /* FIXME: The following is disabled because of PR middle-end/113436.  */
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+#endif
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  #pragma omp parallel if (0) private(A) private(B) private(C) \
+                       allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C)
+    {
+      if (((uintptr_t) &A) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &B) % 128  != 0)
+	abort ();
+      if (((uintptr_t) &C) % 128  != 0)
+	abort ();
+      A = 99;
+      for (int i = 0; i < 10; i++)
+	B[i] = -i-23;
+      C = &A;
+    }
+
+  if (A != 5)
+    abort ();
+  if (c_saved != (uintptr_t) C)
+    abort ();
+  for (int i = 0; i < 10; i++)
+    if (B[i] != i + 5 || C[i] != i+5)
+      abort ();
+
+  free (C);
+}
+
+int
+main ()
+{
+  struct S s;
+  struct St<int> st;
+  for (int dev = 0; dev <= omp_get_num_devices(); dev++)
+    {
+      s.f (dev);
+      st.ft (dev);
+      s.g (dev);
+      st.gt (dev);
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/private-c++-2.C b/libgomp/testsuite/libgomp.c++/private-c++-2.C
new file mode 100644
index 00000000000..aa472cb62ee
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/private-c++-2.C
@@ -0,0 +1,117 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+
+struct t {
+  int A;
+  void f (int dev);
+};
+
+void
+t::f (int dev)
+{
+  int B = 49;
+
+  A = 7;
+  #pragma omp parallel private(A) if(0) shared(B) default(none)
+  {
+    A = 5;
+    B = A;
+  }
+  if (A != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", A); __builtin_abort (); }
+  if (B != 5) { __builtin_printf("ERROR 1a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+  #pragma omp parallel private(A)if(0) shared(B) default(none)
+  {
+    A = 6;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", A); __builtin_abort (); }
+  if (B != 6) { __builtin_printf("ERROR 2a: %d\n", B); __builtin_abort (); }
+  A = 8; B = 49;
+
+  #pragma omp target private(A) map(from:B) device(dev)
+  {
+    A = 7;
+    B = A;
+  }
+  if (A != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", A); __builtin_abort (); }
+  if (B != 7) { __builtin_printf("ERROR 3a: %d\n", B); __builtin_abort (); }
+  A = 9; B = 49;
+  #pragma omp target private(A) map(from:B) device(dev)
+  {
+    A = 8;
+    B = A;
+  }
+  if (A != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", A); __builtin_abort (); }
+  if (B != 8) { __builtin_printf("ERROR 4a: %d\n", B); __builtin_abort (); }
+}
+
+
+template <typename T>
+struct tt {
+  T C;
+  void g (int dev);
+};
+
+template <typename T>
+void
+tt<T>::g (int dev)
+{
+  T D = 49;
+  C = 7;
+  #pragma omp parallel private(C) if(0) shared(D) default(none)
+  {
+    C = 5;
+    D = C;
+  }
+  if (C != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", C);__builtin_abort (); }
+  if (D != 5) { __builtin_printf("ERROR 1a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp parallel private(C)if(0) shared(D) default(none)
+  {
+    C = 6;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", C);__builtin_abort (); }
+  if (D != 6) { __builtin_printf("ERROR 2a: %d\n", D);__builtin_abort (); }
+  C = 8; D = 49;
+  #pragma omp target private(C) map(from:D) defaultmap(none) device(dev)
+  {
+    C = 7;
+    D = C;
+  }
+  if (C != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", C);__builtin_abort (); }
+  if (D != 7) { __builtin_printf("ERROR 3a: %d\n", D);__builtin_abort (); }
+  C = 9; D = 49;
+  #pragma omp target private(C) map(from:D) defaultmap(none) device(dev)
+  {
+    C = 8;
+    D = C;
+  }
+  if (C != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", C); __builtin_abort (); }
+  if (D != 8) { __builtin_printf("ERROR 4a: %d\n", D); }
+}
+
+void
+foo ()
+{
+  struct t x;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    x.f (dev);
+}
+
+void
+bar ()
+{
+  struct tt<int> y;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    y.g (dev);
+}
+
+int
+main ()
+{
+  foo ();
+  bar ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-3.C b/libgomp/testsuite/libgomp.c++/target-lambda-3.C
new file mode 100644
index 00000000000..6be8426bd3e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-3.C
@@ -0,0 +1,104 @@
+// 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>
+#include <omp.h>
+
+template <typename L>
+void
+omp_target_loop (int begin, int end, L loop, int dev)
+{
+  #pragma omp target teams distribute parallel for device(dev)
+  for (int i = begin; i < end; i++)
+    loop (i);
+}
+
+struct S
+{
+  int a, len;
+  int *ptr;
+
+  auto merge_data_func (int *iptr, int &b, int dev)
+  {
+    auto fn = [=](void) -> bool
+      {
+	bool mapped = (omp_target_is_present (iptr, dev)
+                       && omp_target_is_present (ptr, dev));
+	#pragma omp target device(dev)
+	{
+	  if (mapped)
+	    {
+	      for (int i = 0; i < len; i++)
+		ptr[i] += a + b + iptr[i];
+	    }
+	}
+	return mapped;
+      };
+    return fn;
+  }
+};
+
+int x = 1;
+
+void run (int dev)
+{
+  const int N = 10;
+  int *data1 = new int[N];
+  int *data2 = new int[N];
+  memset (data1, 0xab, sizeof (int) * N);
+  memset (data2, 0xcd, sizeof (int) * N);
+
+  bool shared_mem = (omp_target_is_present (data1, dev)
+		     && omp_target_is_present (data2, dev));
+  int val = 1;
+  int &valref = val;
+  #pragma omp target enter data map(alloc: data1[:N], data2[:N]) device(dev)
+
+  omp_target_loop (0, N, [=](int i) { data1[i] = val; }, dev);
+  omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }, dev);
+
+  #pragma omp target update from(data1[:N], data2[:N]) device(dev)
+
+  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]) device(dev)
+
+  int b = 8;
+  S s = { 4, N, data1 };
+  auto f = s.merge_data_func (data2, b, dev);
+  if (f () ^ shared_mem) abort ();
+
+  #pragma omp target enter data map(to: data1[:N]) device(dev)
+  if (f () ^ shared_mem) abort ();
+
+  #pragma omp target enter data map(to: data2[:N]) device(dev)
+  if (!f ()) abort ();
+
+  #pragma omp target exit data map(from: data1[:N], data2[:N]) device(dev)
+
+  for (int i = 0; i < N; i++)
+    {
+      if ((!shared_mem && data1[i] != 0xf)
+	  || (shared_mem && data1[i] != 0x2b))
+	abort ();
+      if (data2[i] != 2) abort ();
+    }
+  delete [] data1;
+  delete [] data2;
+}
+
+int main ()
+{
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    run (dev);
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) firstprivate\(mapped\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(_[0-9]+\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C b/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C
new file mode 100644
index 00000000000..0bb6ce6434b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C
@@ -0,0 +1,125 @@
+/* PR c++/110347 */
+
+#include <omp.h>
+
+#define N 30
+
+struct t {
+  int *A;
+  void f (int dev);
+};
+
+void
+t::f (int dev)
+{
+  int *ptr;
+  int B[N];
+  for (int i = 0; i < N; i++)
+    B[i] = 1 + i;
+  ptr = A = (int *) omp_target_alloc (sizeof (int) * N, dev);
+  omp_target_memcpy (A, B, sizeof (int) * N, 0, 0, dev, omp_initial_device);
+
+  #pragma omp target is_device_ptr (A) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (A[i] != 1 + i)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      A[i] = (-2-i)*10;
+    A = (int *) 0x12345;
+  }
+  if (ptr != A)
+    __builtin_abort ();
+
+  #pragma omp target is_device_ptr (A) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (A[i] != (-2-i)*10)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      A[i] = (3+i)*11;
+    A = (int *) 0x12345;
+  }
+  if (ptr != A)
+    __builtin_abort ();
+
+  int *C = (int *) __builtin_malloc (sizeof(int)*N);
+  omp_target_memcpy (C, A, sizeof (int) * N, 0, 0, omp_initial_device, dev);
+  for (int i = 0; i < N; i++)
+    if (C[i] != (3+i)*11)
+      __builtin_abort ();
+  __builtin_free (C);
+  omp_target_free (A, dev);
+}
+
+template <typename T>
+struct tt {
+  T *D;
+  void g (int dev);
+};
+
+template <typename T>
+void
+tt<T>::g (int dev)
+{
+  T *ptr;
+  T E[N];
+  for (int i = 0; i < N; i++)
+    E[i] = 1 + i;
+  ptr = D = (T *) omp_target_alloc (sizeof (T) * N, dev);
+  omp_target_memcpy (D, E, sizeof (T) * N, 0, 0, dev, omp_initial_device);
+
+  #pragma omp target is_device_ptr (D) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (D[i] != 1 + i)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      D[i] = (-2-i)*10;
+    D = (T *) 0x12345;
+  }
+  if (ptr != D)
+    __builtin_abort ();
+
+  #pragma omp target is_device_ptr (D) device(dev)
+  {
+    for (int i = 0; i < N; i++)
+      if (D[i] != (-2-i)*10)
+	__builtin_abort ();
+    for (int i = 0; i < N; i++)
+      D[i] = (3+i)*11;
+    D = (T *) 0x12345;
+  }
+  if (ptr != D)
+    __builtin_abort ();
+
+  T *F = (T *) __builtin_malloc (sizeof(T)*N);
+  omp_target_memcpy (F, D, sizeof (T) * N, 0, 0, omp_initial_device, dev);
+  for (int i = 0; i < N; i++)
+    if (F[i] != (3+i)*11)
+      __builtin_abort ();
+  __builtin_free (F);
+  omp_target_free (D, dev);}
+
+void
+foo ()
+{
+  struct t x;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    x.f (dev);
+}
+
+void
+bar ()
+{
+  struct tt<int> y;
+  for (int dev = 0; dev <= omp_get_num_devices (); dev++)
+    y.g (dev);
+}
+
+int
+main ()
+{
+  foo ();
+  bar ();
+}
