diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d366703..85b0951 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2696,30 +2696,16 @@
 		  tree present = gfc_omp_check_optional_argument (decl, true);
 		  if (openacc && n->sym->ts.type == BT_CLASS)
 		    {
-		      tree type = TREE_TYPE (decl);
 		      if (n->sym->attr.optional)
 			sorry ("optional class parameter");
-		      if (POINTER_TYPE_P (type))
-			{
-			  node4 = build_omp_clause (input_location,
-						    OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
-			  OMP_CLAUSE_DECL (node4) = decl;
-			  OMP_CLAUSE_SIZE (node4) = size_int (0);
-			  decl = build_fold_indirect_ref (decl);
-			}
 		      tree ptr = gfc_class_data_get (decl);
 		      ptr = build_fold_indirect_ref (ptr);
 		      OMP_CLAUSE_DECL (node) = ptr;
 		      OMP_CLAUSE_SIZE (node) = gfc_class_vtab_size_get (decl);
 		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
-		      OMP_CLAUSE_DECL (node2) = decl;
-		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH);
-		      OMP_CLAUSE_DECL (node3) = gfc_class_data_get (decl);
-		      OMP_CLAUSE_SIZE (node3) = size_int (0);
+		      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_ATTACH_DETACH);
+		      OMP_CLAUSE_DECL (node2) = gfc_class_data_get (decl);
+		      OMP_CLAUSE_SIZE (node2) = size_int (0);
 		      goto finalize_map_clause;
 		    }
 		  else if (POINTER_TYPE_P (TREE_TYPE (decl))
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index bb87128..07baf4d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8323,8 +8323,9 @@
    has array type, else return NULL.  */
 
 static tree
-extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
-			 poly_offset_int *poffsetp, tree *offsetp)
+extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
+			 poly_int64 *bitposp, poly_offset_int *poffsetp,
+			 tree *offsetp, bool openmp)
 {
   tree offset;
   poly_int64 bitsize, bitpos;
@@ -8332,6 +8333,9 @@
   int unsignedp, reversep, volatilep = 0;
   poly_offset_int poffset;
 
+  if (base_ind)
+    *base_ind = NULL_TREE;
+
   if (base_ref)
     *base_ref = NULL_TREE;
 
@@ -8352,14 +8356,29 @@
   base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
 			      &unsignedp, &reversep, &volatilep);
 
-  tree orig_base = base;
-
+  if (!openmp
+      && (TREE_CODE (base) == INDIRECT_REF
+	  || (TREE_CODE (base) == MEM_REF
+	      && integer_zerop (TREE_OPERAND (base, 1))))
+      && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == POINTER_TYPE)
+    {
+      if (base_ind)
+	*base_ind = base;
+      base = TREE_OPERAND (base, 0);
+    }
   if ((TREE_CODE (base) == INDIRECT_REF
        || (TREE_CODE (base) == MEM_REF
 	   && integer_zerop (TREE_OPERAND (base, 1))))
       && DECL_P (TREE_OPERAND (base, 0))
       && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
-    base = TREE_OPERAND (base, 0);
+    {
+      if (base_ref)
+	*base_ref = base;
+      base = TREE_OPERAND (base, 0);
+    }
+
+  if (!openmp)
+    STRIP_NOPS (base);
 
   if (offset && poly_int_tree_p (offset))
     {
@@ -8376,10 +8395,6 @@
   *poffsetp = poffset;
   *offsetp = offset;
 
-  /* Set *BASE_REF if BASE was a dereferenced reference variable.  */
-  if (base_ref && orig_base != base)
-    *base_ref = orig_base;
-
   return base;
 }
 
@@ -8401,6 +8416,48 @@
   return operand_equal_p (expr, base_ptr);
 }
 
+/* Remove COMPONENT_REFS and indirections from EXPR.  */
+
+static tree
+strip_components_and_deref (tree expr)
+{
+  while (TREE_CODE (expr) == COMPONENT_REF
+	 || TREE_CODE (expr) == INDIRECT_REF
+	 || (TREE_CODE (expr) == MEM_REF
+	     && integer_zerop (TREE_OPERAND (expr, 1))))
+    expr = TREE_OPERAND (expr, 0);
+
+  return expr;
+}
+
+/* Return TRUE if EXPR is something we will use as the base of an aggregate
+   access, either:
+
+  - a DECL_P.
+  - a struct component with no indirection ("a.b.c").
+  - a struct component with indirection ("a->b->c").
+*/
+
+static bool
+aggregate_base_p (tree expr)
+{
+  while (TREE_CODE (expr) == COMPONENT_REF
+	 && (DECL_P (TREE_OPERAND (expr, 0))
+	     || (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF)))
+    expr = TREE_OPERAND (expr, 0);
+
+  if (DECL_P (expr))
+    return true;
+
+  if (TREE_CODE (expr) == COMPONENT_REF
+      && (TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF
+	  || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF
+	      && integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1)))))
+    return true;
+
+  return false;
+}
+
 /* Implement OpenMP 5.x map ordering rules for target directives. There are
    several rules, and with some level of ambiguity, hopefully we can at least
    collect the complexity here in one place.  */
@@ -8788,10 +8845,12 @@
 {
   poly_offset_int coffset;
   poly_int64 cbitpos;
-  tree base_ref, tree_coffset;
+  tree base_ind, base_ref, tree_coffset;
+  bool openmp = !(region_type & ORT_ACC);
 
-  tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
-				       &cbitpos, &coffset, &tree_coffset);
+  tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ind,
+				       &base_ref, &cbitpos, &coffset,
+				       &tree_coffset, openmp);
 
   bool do_map_struct = (base == decl && !tree_coffset);
 
@@ -8830,12 +8889,7 @@
     return NULL_TREE;
 
   /* Nor for attach_detach for OpenMP.  */
-  if ((code == OMP_TARGET
-       || code == OMP_TARGET_DATA
-       || code == OMP_TARGET_UPDATE
-       || code == OMP_TARGET_ENTER_DATA
-       || code == OMP_TARGET_EXIT_DATA)
-      && attach_detach)
+  if (openmp && attach_detach)
     {
       if (DECL_P (decl))
 	{
@@ -8858,12 +8912,15 @@
 
       OMP_CLAUSE_SET_MAP_KIND (l, k);
 
-      if (base_ref)
+      if (!openmp && base_ind)
+	OMP_CLAUSE_DECL (l) = unshare_expr (base_ind);
+      else if (base_ref)
 	OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
       else
 	{
 	  OMP_CLAUSE_DECL (l) = unshare_expr (decl);
-	  if (!DECL_P (OMP_CLAUSE_DECL (l))
+	  if (openmp
+	      && !DECL_P (OMP_CLAUSE_DECL (l))
 	      && (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL,
 				 is_gimple_lvalue, fb_lvalue) == GS_ERROR))
 	    return error_mark_node;
@@ -8919,6 +8976,48 @@
 	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
 	  OMP_CLAUSE_CHAIN (l) = c2;
 	}
+      else if (!openmp
+	       && (base_ind || base_ref)
+	       && (region_type & ORT_TARGET))
+	{
+	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+	  enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
+					      : GOMP_MAP_FIRSTPRIVATE_POINTER;
+	  OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+	  OMP_CLAUSE_SIZE (c2) = size_zero_node;
+	  tree sdecl = strip_components_and_deref (decl);
+	  if (DECL_P (decl)
+	      && (POINTER_TYPE_P (TREE_TYPE (sdecl))
+		  || TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE))
+	    {
+	      /* Insert after struct node.  */
+	      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
+	      OMP_CLAUSE_DECL (c2) = decl;
+	      OMP_CLAUSE_CHAIN (l) = c2;
+	    }
+	  else
+	    {
+	      /* If the ultimate base for this component access is not a
+		 pointer or reference, that means it is a struct component
+		 access itself.  Insert a node to be processed on the next
+		 iteration of our caller's loop, which will subsequently be
+		 turned into a new GOMP_MAP_STRUCT mapping itself.
+
+		 We need to do this else the non-DECL_P base won't be
+		 rewritten correctly in the offloaded region.  */
+	      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					  OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
+	      OMP_CLAUSE_DECL (c2) = unshare_expr (decl);
+	      OMP_CLAUSE_SIZE (c2) = (DECL_P (decl)
+				      ? DECL_SIZE_UNIT (decl)
+				      : TYPE_SIZE_UNIT (TREE_TYPE (decl)));
+	      tree *next_node = &OMP_CLAUSE_CHAIN (*list_p);
+	      OMP_CLAUSE_CHAIN (c2) = *next_node;
+	      *next_node = c2;
+	      return NULL_TREE;
+	    }
+	}
       *flags = GOVD_MAP | GOVD_EXPLICIT;
       if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
 	*flags |= GOVD_SEEN;
@@ -8927,7 +9026,8 @@
 
       /* If this is a *pointer-to-struct expression, make sure a
 	 firstprivate map of the base-pointer exists.  */
-      if (component_ref_p
+      if (openmp
+	  && component_ref_p
 	  && ((TREE_CODE (decl) == MEM_REF
 	       && integer_zerop (TREE_OPERAND (decl, 1)))
 	      || INDIRECT_REF_P (decl))
@@ -8957,10 +9057,12 @@
 	n->value |= GOVD_SEEN;
       sc = &OMP_CLAUSE_CHAIN (*osc);
       /* The struct mapping might be immediately followed by a
-	 FIRSTPRIVATE_REFERENCE if it is a reference.  (This added node is
-	 removed in omp-low.c after it has been processed there.)  */
+	 FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an
+	 indirect access or a reference, or both.  (This added node is removed
+	 in omp-low.c after it has been processed there.)  */
       if (*sc != c
-	  && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	  && (OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	      || OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 	sc = &OMP_CLAUSE_CHAIN (*sc);
       for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
 	if ((ptr || attach_detach) && sc == prev_list_p)
@@ -8975,9 +9077,10 @@
 	    poly_offset_int offset;
 	    poly_int64 bitpos;
 	    tree tree_offset;
-	    tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos,
-						 &offset, &tree_offset);
-	    if (base != decl)
+	    tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
+						 &bitpos, &offset,
+						 &tree_offset, openmp);
+	    if (!base || !operand_equal_p (base, decl, 0))
 	      break;
 	    if (scp)
 	      continue;
@@ -9098,8 +9201,9 @@
 	}
       else if (*sc != c)
 	{
-	  if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
-	      == GS_ERROR)
+	  if (openmp
+	      && (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+		  == GS_ERROR))
 	    return error_mark_node;
 	  /* In the non-pointer case, the mapping clause itself is moved into
 	     the correct position in the struct component list, which in this
@@ -9702,10 +9806,43 @@
 	      tree indir_base = NULL_TREE;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
-	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
-		  && TREE_CODE (*pd) == COMPONENT_REF
-		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
-		  && code != OACC_UPDATE)
+	      if ((region_type & ORT_ACC) && TREE_CODE (decl) == COMPONENT_REF)
+		{
+		  /* Strip off component refs from RHS of e.g. "a->b->c.d.e"
+		     (which would leave "a->b" in that case).  This is intended
+		     to be equivalent to the base finding done by
+		     get_inner_reference.  */
+		  while (TREE_CODE (decl) == COMPONENT_REF
+			 && (DECL_P (TREE_OPERAND (decl, 0))
+			     || (TREE_CODE (TREE_OPERAND (decl, 0))
+				 == COMPONENT_REF)))
+		    decl = TREE_OPERAND (decl, 0);
+
+		  if (TREE_CODE (decl) == COMPONENT_REF)
+		    decl = TREE_OPERAND (decl, 0);
+
+		  /* Strip off RHS from "a->b".  */
+		  if ((TREE_CODE (decl) == INDIRECT_REF
+		       || (TREE_CODE (decl) == MEM_REF
+			   && integer_zerop (TREE_OPERAND (decl, 1))))
+		      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			  == POINTER_TYPE))
+		    decl = TREE_OPERAND (decl, 0);
+
+		  /* Strip off RHS from "a_ref.b" (where a_ref is
+		     reference-typed).  */
+		  if (TREE_CODE (decl) == INDIRECT_REF
+		      && DECL_P (TREE_OPERAND (decl, 0))
+		      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			  == REFERENCE_TYPE))
+		    decl = TREE_OPERAND (decl, 0);
+
+		  STRIP_NOPS (decl);
+		}
+	      else if ((region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+		       && TREE_CODE (*pd) == COMPONENT_REF
+		       && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+		       && code != OACC_UPDATE)
 		{
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    {
@@ -9808,11 +9945,13 @@
 	      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)
-		   || (component_ref_p
-		       && (INDIRECT_REF_P (decl)
-			   || TREE_CODE (decl) == MEM_REF
-			   || TREE_CODE (decl) == ARRAY_REF)))
+	      if ((((region_type & ORT_ACC) && aggregate_base_p (decl))
+		   || (!(region_type & ORT_ACC)
+		       && (DECL_P (decl)
+			   || (component_ref_p
+			       && (INDIRECT_REF_P (decl)
+				   || TREE_CODE (decl) == MEM_REF
+				   || TREE_CODE (decl) == ARRAY_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
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 64b7c19..bb473d5 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1651,8 +1651,10 @@
 	      if (TREE_CODE (decl) == COMPONENT_REF
 		  || (TREE_CODE (decl) == INDIRECT_REF
 		      && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
-		      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
-			  == REFERENCE_TYPE)))
+		      && (((TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			    == REFERENCE_TYPE)
+			   || (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			       == POINTER_TYPE)))))
 		break;
 	      if (DECL_SIZE (decl)
 		  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -13504,6 +13506,7 @@
 		if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
 		  is_ref = false;
 		bool ref_to_array = false;
+		bool ref_to_ptr = false;
 		if (is_ref)
 		  {
 		    type = TREE_TYPE (type);
@@ -13522,6 +13525,12 @@
 		    new_var = decl2;
 		    type = TREE_TYPE (new_var);
 		  }
+		else if (TREE_CODE (type) == REFERENCE_TYPE
+			 && TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE)
+		  {
+		    type = TREE_TYPE (type);
+		    ref_to_ptr = true;
+		  }
 		x = build_receiver_ref (prev, false, ctx);
 		x = fold_convert_loc (clause_loc, type, x);
 		if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
@@ -13544,7 +13553,8 @@
 		if (ref_to_array)
 		  x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
 		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
-		if (is_ref && !ref_to_array)
+		if ((is_ref && !ref_to_array)
+		    || ref_to_ptr)
 		  {
 		    tree t = create_tmp_var_raw (type, get_name (var));
 		    gimple_add_tmp_var (t);
diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
new file mode 100644
index 0000000..e0c1157
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
@@ -0,0 +1,13 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct Foo {
+  float *a;
+  void init(int N) {
+    a = new float[N];
+    #pragma acc enter data create(a[0:N])
+  }
+};
+int main() { Foo x; x.init(1024); }
+
+/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/member-array-omp.C b/gcc/testsuite/g++.dg/gomp/member-array-omp.C
new file mode 100644
index 0000000..a53aa44
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/member-array-omp.C
@@ -0,0 +1,13 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct Foo {
+  float *a;
+  void init(int N) {
+    a = new float[N];
+    #pragma omp target enter data map(alloc:a[0:N])
+  }
+};
+int main() { Foo x; x.init(1024); }
+
+/* { dg-final { scan-tree-dump {map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:this->a \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
new file mode 100644
index 0000000..dacbb52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
@@ -0,0 +1,101 @@
+#include <cassert>
+
+/* Test attach/detach operation with pointers and references to structs.  */
+
+typedef struct mystruct {
+  int *a;
+  int b;
+  int *c;
+  int d;
+  int *e;
+} mystruct;
+
+void str (void)
+{
+  int a[10], c[10], e[10];
+  mystruct m = { .a = a, .c = c, .e = e };
+  a[0] = 5;
+  c[0] = 7;
+  e[0] = 9;
+  #pragma acc parallel copy(m.a[0:10], m.b, m.c[0:10], m.d, m.e[0:10])
+  {
+    m.a[0] = m.c[0] + m.e[0];
+  }
+  assert (m.a[0] == 7 + 9);
+}
+
+void strp (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct *m = new mystruct;
+  m->a = a;
+  m->c = c;
+  m->e = e;
+  a[0] = 6;
+  c[0] = 8;
+  e[0] = 10;
+  #pragma acc parallel copy(m->a[0:10], m->b, m->c[0:10], m->d, m->e[0:10])
+  {
+    m->a[0] = m->c[0] + m->e[0];
+  }
+  assert (m->a[0] == 8 + 10);
+  delete m;
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strr (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct m;
+  mystruct &n = m;
+  n.a = a;
+  n.c = c;
+  n.e = e;
+  a[0] = 7;
+  c[0] = 9;
+  e[0] = 11;
+  #pragma acc parallel copy(n.a[0:10], n.b, n.c[0:10], n.d, n.e[0:10])
+  {
+    n.a[0] = n.c[0] + n.e[0];
+  }
+  assert (n.a[0] == 9 + 11);
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strrp (void)
+{
+  int a[10], c[10], e[10];
+  mystruct *m = new mystruct;
+  mystruct *&n = m;
+  n->a = a;
+  n->b = 3;
+  n->c = c;
+  n->d = 5;
+  n->e = e;
+  a[0] = 8;
+  c[0] = 10;
+  e[0] = 12;
+  #pragma acc parallel copy(n->a[0:10], n->c[0:10], n->e[0:10])
+  {
+    n->a[0] = n->c[0] + n->e[0];
+  }
+  assert (n->a[0] == 10 + 12);
+  delete m;
+}
+
+int main (int argc, char *argv[])
+{
+  str ();
+  strp ();
+  strr ();
+  strrp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
new file mode 100644
index 0000000..27fe1a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
new file mode 100644
index 0000000..a7308e8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,231 @@
+#include <stdlib.h>
+
+/* Test mapping chained indirect struct accesses, mixed in different ways.  */
+
+typedef struct {
+  int *a;
+  int b;
+  int *c;
+} str1;
+
+typedef struct {
+  int d;
+  int *e;
+  str1 *f;
+} str2;
+
+typedef struct {
+  int g;
+  int h;
+  str2 *s2;
+} str3;
+
+typedef struct {
+  str3 m;
+  str3 n;
+} str4;
+
+void
+zero_arrays (str4 *s, int N)
+{
+  for (int i = 0; i < N; i++)
+    {
+      s->m.s2->e[i] = 0;
+      s->m.s2->f->a[i] = 0;
+      s->m.s2->f->c[i] = 0;
+      s->n.s2->e[i] = 0;
+      s->n.s2->f->a[i] = 0;
+      s->n.s2->f->c[i] = 0;
+    }
+}
+
+void
+alloc_s2 (str2 **s, int N)
+{
+  (*s) = (str2 *) malloc (sizeof (str2));
+  (*s)->f = (str1 *) malloc (sizeof (str1));
+  (*s)->e = (int *) malloc (sizeof (int) * N);
+  (*s)->f->a = (int *) malloc (sizeof (int) * N);
+  (*s)->f->c = (int *) malloc (sizeof (int) * N);
+}
+
+int main (int argc, char* argv[])
+{
+  const int N = 1024;
+  str4 p, *q;
+  int i;
+
+  alloc_s2 (&p.m.s2, N);
+  alloc_s2 (&p.n.s2, N);
+  q = (str4 *) malloc (sizeof (str4));
+  alloc_s2 (&q->m.s2, N);
+  alloc_s2 (&q->n.s2, N);
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	p.m.s2->e[j]++;
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
+			  copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
+	|| p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
+		       copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
+		      copyout(p.n.s2->e[:N])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
+	|| p.n.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc parallel loop copy(q->m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	q->m.s2->e[j]++;
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
+			  copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
+	|| q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
+		       copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
+		      copyout(q->n.s2->e[:N])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
+	|| q->n.s2->e[i] != 99)
+      abort ();
+
+  return 0;
+}