diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index e3df4bbf84e..9098b35c9f1 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2693,30 +2693,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  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 86000f8470b..6d204908c82 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8323,8 +8323,8 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
    has array type, else return NULL.  */
 
 static tree
-extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
-			 poly_offset_int *poffsetp)
+extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
+			 poly_int64 *bitposp, poly_offset_int *poffsetp)
 {
   tree offset;
   poly_int64 bitsize, bitpos;
@@ -8332,6 +8332,9 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   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 +8355,27 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
 			      &unsignedp, &reversep, &volatilep);
 
-  tree orig_base = base;
-
+  if ((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);
+    }
+
+  STRIP_NOPS (base);
 
   gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset));
 
@@ -8374,10 +8390,6 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   *bitposp = bitpos;
   *poffsetp = poffset;
 
-  /* Set *BASE_REF if BASE was a dereferenced reference variable.  */
-  if (base_ref && orig_base != base)
-    *base_ref = orig_base;
-
   return base;
 }
 
@@ -8394,6 +8406,59 @@ is_or_contains_p (tree expr, tree base_ptr)
   return expr == base_ptr;
 }
 
+/* Remove COMPONENT_REFS from EXPR.  */
+
+static tree
+strip_components (tree expr)
+{
+  while (TREE_CODE (expr) == COMPONENT_REF)
+    expr = TREE_OPERAND (expr, 0);
+
+  return expr;
+}
+
+/* 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.  */
@@ -8686,20 +8751,24 @@ move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr,
 static tree
 build_struct_group (struct gimplify_omp_ctx *ctx,
 		    enum omp_region_type region_type, enum tree_code code,
-		    tree decl, tree *pd, unsigned int *flags, tree c,
+		    tree decl, unsigned int *flags, tree c,
 		    hash_map<tree_operand_hash, tree> *&struct_map_to_clause,
 		    tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p,
 		    bool *cont)
 {
   poly_offset_int coffset;
   poly_int64 cbitpos;
-  tree base_ref;
+  tree base_ind, base_ref;
+  tree *list_in_p = list_p, *prev_list_in_p = prev_list_p;
 
-  tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
-				       &cbitpos, &coffset);
+  tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ind,
+				       &base_ref, &cbitpos, &coffset);
 
   gcc_assert (base == decl);
 
+  /* Here, DECL is usually a DECL_P, unless we have chained indirect member
+     accesses, e.g. mystruct->a->b.  In that case it'll be the "mystruct->a"
+     part.  */
   splay_tree_node n
     = (DECL_P (decl)
        ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl)
@@ -8733,7 +8802,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
       gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT;
 
       OMP_CLAUSE_SET_MAP_KIND (l, k);
-      if (base_ref)
+      if (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) = decl;
@@ -8763,15 +8834,81 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
 	}
       else
 	list_p = insert_node_after (l, list_p);
-      if (base_ref && code == OMP_TARGET)
+
+	/* Handle pointers to structs and references to structs: these cases
+	   have an additional GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} node
+	   inserted after the GOMP_MAP_STRUCT node.  References to pointers
+	   use GOMP_MAP_FIRSTPRIVATE_REFERENCE.  */
+      if ((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 = GOMP_MAP_FIRSTPRIVATE_REFERENCE;
+	  enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
+					      : 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 (l);
-	  OMP_CLAUSE_CHAIN (l) = c2;
+	  tree sdecl = strip_components (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 (!POINTER_TYPE_P (TREE_TYPE (sdecl))
+		   && TREE_CODE (TREE_TYPE (sdecl)) != REFERENCE_TYPE)
+	    {
+	      /* 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;
+	      *cont = true;
+	      return NULL_TREE;
+	    }
+	  else
+	    {
+	      /* We have chained indirect structure member accesses, e.g.
+		 "mystruct->a->b".  DECL isn't a DECL_P in that case (it will
+		 be "mystruct->a"), and we can't use
+		 GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} directly after the
+		 GOMP_MAP_STRUCT.  Create a new FORCE_PRESENT mapping instead
+		 for the ultimate root node (i.e. "mystruct"), and use
+		 that.  */
+	      tree use_decl = unshare_expr (sdecl);
+	      decl = strip_components_and_deref (decl);
+	      OMP_CLAUSE_DECL (c2) = decl;
+	      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					  OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_FORCE_PRESENT);
+	      OMP_CLAUSE_SIZE (c3) = size_one_node;
+	      OMP_CLAUSE_DECL (c3) = use_decl;
+	      /* Insert FORCE_PRESENT and FIRSTPRIVATE_{REFERENCE,POINTER}
+		 nodes before struct node.  */
+	      OMP_CLAUSE_CHAIN (c3) = c2;
+	      /* If we have a pointer mapping group, insert these nodes before
+		 the start of that group, else insert before incoming list_p
+		 (the currently-scanned node).  */
+	      tree *insert_at
+		= (ptr || attach_detach) ? prev_list_in_p : list_in_p;
+	      OMP_CLAUSE_CHAIN (c2) = *insert_at;
+	      *insert_at = c3;
+	      list_p = &OMP_CLAUSE_CHAIN (c2);
+	    }
 	}
       *flags = GOVD_MAP | GOVD_EXPLICIT;
       if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
@@ -8784,11 +8921,18 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
     {
       tree *osc = struct_map_to_clause->get (decl);
       tree *sc = NULL, *scp = NULL;
-      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
+      if (n && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+		|| ptr
+		|| attach_detach))
 	n->value |= GOVD_SEEN;
       sc = &OMP_CLAUSE_CHAIN (*osc);
+      /* The struct mapping might be immediately followed by a
+	 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.)  */
       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))
 	{
@@ -8803,8 +8947,8 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
 	      tree sc_decl = OMP_CLAUSE_DECL (*sc);
 	      poly_offset_int offset;
 	      poly_int64 bitpos;
-	      tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos,
-						   &offset);
+	      tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
+						   &bitpos, &offset);
 	      if (base != decl)
 		break;
 	      if (scp)
@@ -8927,9 +9071,6 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
 	}
       else if (*sc != c)
 	{
-	  if (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
 	     case is just SC.  */
@@ -8951,7 +9092,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
   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;
   int nowait = -1;
@@ -9518,111 +9658,43 @@ 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)
-		  && (OMP_CLAUSE_MAP_KIND (c)
-		      != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
+		      == REFERENCE_TYPE))
 		{
 		  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
-		  && TREE_CODE (*pd) == COMPONENT_REF
-		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
-		  && code != OACC_UPDATE)
+	      if (TREE_CODE (decl) == COMPONENT_REF)
 		{
-		  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))
-			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
-			      == POINTER_TYPE))
-			{
-			  indir_p = true;
-			  decl = TREE_OPERAND (decl, 0);
-			}
-		      if (TREE_CODE (decl) == INDIRECT_REF
-			  && DECL_P (TREE_OPERAND (decl, 0))
-			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
-			      == REFERENCE_TYPE))
-			{
-			  decl_ref = decl;
-			  decl = TREE_OPERAND (decl, 0);
-			}
-		    }
-		}
-	      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)
+		  /* 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);
-		}
-	      if (decl != orig_decl && DECL_P (decl) && indir_p)
-		{
-		  gomp_map_kind k
-		    = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
-		       ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
-		  /* We have a dereference of a struct member.  Make this an
-		     attach/detach operation, and ensure the base pointer is
-		     mapped as a FIRSTPRIVATE_POINTER.  */
-		  OMP_CLAUSE_SET_MAP_KIND (c, k);
-		  flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
-		  tree next_clause = OMP_CLAUSE_CHAIN (c);
-		  if (k == GOMP_MAP_ATTACH
-		      && code != OACC_ENTER_DATA
-		      && code != OMP_TARGET_ENTER_DATA
-		      && (!next_clause
-			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
-			   || (OMP_CLAUSE_MAP_KIND (next_clause)
-			       != GOMP_MAP_POINTER)
-			   || OMP_CLAUSE_DECL (next_clause) != decl)
-		      && (!struct_deref_set
-			  || !struct_deref_set->contains (decl)))
-		    {
-		      if (!struct_deref_set)
-			struct_deref_set = new hash_set<tree> ();
-		      /* As well as the attach, we also need a
-			 FIRSTPRIVATE_POINTER clause to properly map the
-			 pointer to the struct base.  */
-		      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-						  OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
-		      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
-			= 1;
-		      tree charptr_zero
-			= build_int_cst (build_pointer_type (char_type_node),
-					 0);
-		      OMP_CLAUSE_DECL (c2)
-			= build2 (MEM_REF, char_type_node,
-				  decl_ref ? decl_ref : decl, charptr_zero);
-		      OMP_CLAUSE_SIZE (c2) = size_zero_node;
-		      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-						  OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (c3,
-					       GOMP_MAP_FIRSTPRIVATE_POINTER);
-		      OMP_CLAUSE_DECL (c3) = decl;
-		      OMP_CLAUSE_SIZE (c3) = size_zero_node;
-		      tree mapgrp = *prev_list_p;
-		      *prev_list_p = c2;
-		      OMP_CLAUSE_CHAIN (c3) = mapgrp;
-		      OMP_CLAUSE_CHAIN (c2) = c3;
-
-		      struct_deref_set->add (decl);
-		    }
-		  goto do_add_decl;
+
+		  STRIP_NOPS (decl);
 		}
 	      /* An "attach/detach" operation on an update directive should
 		 behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that
@@ -9631,8 +9703,7 @@ 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)
-		   || (component_ref_p && INDIRECT_REF_P (decl)))
+	      if (aggregate_base_p (decl)
 		  && 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
@@ -9679,7 +9750,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		    }
 		  bool cont = false;
 		  tree add_decl
-		    = build_struct_group (ctx, region_type, code, decl, pd,
+		    = build_struct_group (ctx, region_type, code, decl,
 					  &flags, c, struct_map_to_clause,
 					  prev_list_p, list_p, pre_p, &cont);
 		  if (add_decl == error_mark_node)
@@ -10310,8 +10381,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   gimplify_omp_ctxp = ctx;
   if (struct_map_to_clause)
     delete struct_map_to_clause;
-  if (struct_deref_set)
-    delete struct_deref_set;
 }
 
 /* Return true if DECL is a candidate for shared to firstprivate
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 64b7c19ff39..bb473d58f12 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1651,8 +1651,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	      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 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		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 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    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 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		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 00000000000..e0c11570f5d
--- /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 00000000000..40fc503f882
--- /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 {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/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C
index f4d40ec8e4b..9fc0e4173d2 100644
--- a/gcc/testsuite/g++.dg/gomp/target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -15,6 +15,8 @@ S::bar (int x)
   #pragma omp target enter data map (alloc: a, b)
 }
 
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct S \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:\\(\\(struct S \\*\\) this\\)->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(\\(struct S \\*\\) this\\)->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+
 template <int N>
 struct T
 {
@@ -33,4 +35,4 @@ T<N>::bar (int x)
 
 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" } } */
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:\\(\\(struct T \\*\\) this\\)->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(\\(struct T \\*\\) this\\)->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index 679c85a54dd..a5e832130fb 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -46,4 +46,4 @@ int main (void)
   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" } } */
+/* { dg-final { scan-tree-dump {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\]\)} "gimple" } } */
diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c
index 08e42eeb304..3e7921270c9 100644
--- a/gcc/testsuite/gcc.dg/gomp/target-3.c
+++ b/gcc/testsuite/gcc.dg/gomp/target-3.c
@@ -13,4 +13,4 @@ void foo (struct S *s)
   #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" { xfail *-*-* } } } */
+/* { 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/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 00000000000..27fe1a9d07d
--- /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 00000000000..f62e190655d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,95 @@
+#include <stdlib.h>
+
+/* Test sorting of nested struct accesses (unfinished!).  */
+
+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;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  str4 p, *q;
+  int i;
+
+  p.m.s2 = (str2 *) malloc (sizeof (str2));
+  p.m.s2->f = (str1 *) malloc (sizeof (str1));
+  p.m.s2->e = (int *) malloc (sizeof (int) * N);
+  p.m.s2->f->a = (int *) malloc (sizeof (int) * N);
+  p.m.s2->f->c = (int *) malloc (sizeof (int) * N);
+
+  p.n.s2 = (str2 *) malloc (sizeof (str2));
+  p.n.s2->f = (str1 *) malloc (sizeof (str1));
+  p.n.s2->e = (int *) malloc (sizeof (int) * N);
+  p.n.s2->f->a = (int *) malloc (sizeof (int) * N);
+  p.n.s2->f->c = (int *) malloc (sizeof (int) * N);
+
+  q = (str4 *) malloc (sizeof (str4));
+
+  q->m.s2 = (str2 *) malloc (sizeof (str2));
+  q->m.s2->f = (str1 *) malloc (sizeof (str1));
+  q->m.s2->e = (int *) malloc (sizeof (int) * N);
+  q->m.s2->f->a = (int *) malloc (sizeof (int) * N);
+  q->m.s2->f->c = (int *) malloc (sizeof (int) * N);
+
+  q->n.s2 = (str2 *) malloc (sizeof (str2));
+  q->n.s2->f = (str1 *) malloc (sizeof (str1));
+  q->n.s2->e = (int *) malloc (sizeof (int) * N);
+  q->n.s2->f->a = (int *) malloc (sizeof (int) * N);
+  q->n.s2->f->c = (int *) malloc (sizeof (int) * N);
+
+  for (i = 0; i < N; i++)
+    {
+      p.m.s2->e[i] = 0;
+      p.m.s2->f->a[i] = 0;
+      p.m.s2->f->c[i] = 0;
+      p.n.s2->e[i] = 0;
+      p.n.s2->f->a[i] = 0;
+      p.n.s2->f->c[i] = 0;
+      q->m.s2->e[i] = 0;
+      q->m.s2->f->a[i] = 0;
+      q->m.s2->f->c[i] = 0;
+      q->n.s2->e[i] = 0;
+      q->n.s2->f->a[i] = 0;
+      q->n.s2->f->c[i] = 0;
+    }
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc enter data copyin(p.m.s2[:1])
+      {
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+	for (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 ();
+    }
+
+  return 0;
+}
