[05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification

Message ID 00411b16d398cb33a52f4357a527adf4bca4a94d.1633106213.git.julian@codesourcery.com
State New
Headers
Series OpenMP: Deep struct dereferences |

Commit Message

Julian Brown Oct. 1, 2021, 5:09 p.m. UTC
  This patch lifts struct sibling-list handling out of the main loop in
gimplify_scan_omp_clauses.  The reasons for this are several: first,
it means that we can subject created sibling list groups to topological
sorting (see previous patch) so base-pointer data dependencies are
handled correctly.

Secondly, it means that in the first pass gathering up sibling lists
from parsed OpenMP/OpenACC clauses, we don't need to worry about
gimplifying: that means we can see struct bases & components we need
to sort sibling lists properly, even when we're using a non-DECL_P
struct base.  Gimplification proper still happens

Thirdly, because we use more than one pass through the clause list and
gather appropriate data, we can tell if we're mapping a whole struct
in a different node, and avoid building struct sibling lists for that
struct appropriately.

Fourthly, we can re-use the node grouping functions from the
previous patch, and thus mostly avoid the "prev_list_p" handling in
gimplify_scan_omp_clauses that tracks the first node in such groups
at present.

Some redundant code has been removed and code paths for OpenACC/OpenMP are
now shared where appropriate, though OpenACC doesn't do the topological
sorting of nodes (yet?).

OK for mainline?

Thanks,

Julian

2021-09-29  Julian Brown  <julian@codesourcery.com>

gcc/
	* gimplify.c (gimplify_omp_var_data): Remove GOVD_MAP_HAS_ATTACHMENTS.
	(extract_base_bit_offset): Remove OFFSETP parameter.
	(strip_components_and_deref): Extend with POINTER_PLUS_EXPR and
	COMPOUND_EXPR handling.
	(aggregate_base_p): Remove.
	(omp_group_last, omp_group_base): Add GOMP_MAP_STRUCT handling.
	(build_struct_group): Remove CTX, DECL, PD, COMPONENT_REF_P, FLAGS,
	STRUCT_SEEN_CLAUSE, PRE_P, CONT parameters.  Replace PREV_LIST_P and C
	parameters with GRP_START_P and GRP_END.  Add INNER.  Update calls to
	extract_base_bit_offset.  Remove gimplification of clauses for OpenMP.
	Rework inner struct handling for OpenACC.  Don't use context's
	variables splay tree.
	(omp_build_struct_sibling_lists): New function, extracted from
	gimplify_scan_omp_clauses and refactored.
	(gimplify_scan_omp_clauses): Call above function to handle struct
	sibling lists.  Remove STRUCT_MAP_TO_CLAUSE, STRUCT_SEEN_CLAUSE,
	STRUCT_DEREF_SET.  Rework flag handling, adding decl for struct
	variables.
	(gimplify_adjust_omp_clauses_1): Remove GOVD_MAP_HAS_ATTACHMENTS
	handling, unused now.

gcc/testsuite/
	* g++.dg/goacc/member-array-acc.C: Update expected output.
	* g++.dg/gomp/target-3.C: Likewise.
	* g++.dg/gomp/target-lambda-1.C: Likewise.
	* g++.dg/gomp/target-this-2.C: Likewise.
	* g++.dg/gomp/target-this-4.C: Likewise.
---
 gcc/gimplify.c                                | 943 ++++++++----------
 gcc/testsuite/g++.dg/goacc/member-array-acc.C |   2 +-
 gcc/testsuite/g++.dg/gomp/target-3.C          |   4 +-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C   |   2 +-
 gcc/testsuite/g++.dg/gomp/target-this-2.C     |   2 +-
 gcc/testsuite/g++.dg/gomp/target-this-4.C     |   4 +-
 6 files changed, 410 insertions(+), 547 deletions(-)
  

Comments

Julian Brown Oct. 1, 2021, 5:16 p.m. UTC | #1
Oops, editing:

On Fri, 1 Oct 2021 10:09:03 -0700
Julian Brown <julian@codesourcery.com> wrote:

> Secondly, it means that in the first pass gathering up sibling lists
> from parsed OpenMP/OpenACC clauses, we don't need to worry about
> gimplifying: that means we can see struct bases & components we need
> to sort sibling lists properly, even when we're using a non-DECL_P
> struct base.  Gimplification proper still happens

...in the main loop in gimplify_scan_omp_clauses.
  

Patch

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index c10a3e8842a..31e2e4d9fe7 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -125,10 +125,6 @@  enum gimplify_omp_var_data
   /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause.  */
   GOVD_REDUCTION_INSCAN = 0x2000000,
 
-  /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
-     fields.  */
-  GOVD_MAP_HAS_ATTACHMENTS = 0x4000000,
-
   /* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT.  */
   GOVD_FIRSTPRIVATE_IMPLICIT = 0x8000000,
 
@@ -8642,7 +8638,7 @@  build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
 
 static tree
 extract_base_bit_offset (tree base, poly_int64 *bitposp,
-			 poly_offset_int *poffsetp, tree *offsetp)
+			 poly_offset_int *poffsetp)
 {
   tree offset;
   poly_int64 bitsize, bitpos;
@@ -8670,7 +8666,6 @@  extract_base_bit_offset (tree base, poly_int64 *bitposp,
 
   *bitposp = bitpos;
   *poffsetp = poffset;
-  *offsetp = offset;
 
   return base;
 }
@@ -8683,8 +8678,15 @@  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);
+	     && integer_zerop (TREE_OPERAND (expr, 1)))
+	 || TREE_CODE (expr) == POINTER_PLUS_EXPR
+	 || TREE_CODE (expr) == COMPOUND_EXPR)
+      if (TREE_CODE (expr) == COMPOUND_EXPR)
+	expr = TREE_OPERAND (expr, 1);
+      else
+	expr = TREE_OPERAND (expr, 0);
+
+  STRIP_NOPS (expr);
 
   return expr;
 }
@@ -8700,34 +8702,6 @@  strip_indirections (tree expr)
   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;
-}
-
 enum omp_tsort_mark {
   UNVISITED,
   TEMPORARY,
@@ -8956,6 +8930,18 @@  omp_group_last (tree *start_p)
 	      || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH))
 	grp_last_p = &OMP_CLAUSE_CHAIN (c);
       break;
+
+    case GOMP_MAP_STRUCT:
+      {
+	unsigned HOST_WIDE_INT num_mappings
+	  = tree_to_uhwi (OMP_CLAUSE_SIZE (c));
+	if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	  grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p);
+	for (unsigned i = 0; i < num_mappings; i++)
+	  grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p);
+      }
+      break;
     }
 
   return grp_last_p;
@@ -9089,6 +9075,21 @@  omp_group_base (omp_mapping_group *grp, unsigned int *chained,
 	internal_error ("unexpected mapping node");
       return error_mark_node;
 
+    case GOMP_MAP_STRUCT:
+      {
+	unsigned HOST_WIDE_INT num_mappings
+	  = tree_to_uhwi (OMP_CLAUSE_SIZE (node));
+	node = OMP_CLAUSE_CHAIN (node);
+	if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	    || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	  {
+	    *firstprivate = OMP_CLAUSE_DECL (node);
+	    node = OMP_CLAUSE_CHAIN (node);
+	  }
+	*chained = num_mappings;
+	return node;
+      }
+
     case GOMP_MAP_FORCE_DEVICEPTR:
     case GOMP_MAP_DEVICE_RESIDENT:
     case GOMP_MAP_LINK:
@@ -9751,21 +9752,16 @@  move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr,
    next node.  PREV_LIST_P and LIST_P may be modified by the function when a
    list rearrangement has taken place.  */
 
-static tree
-build_struct_group (struct gimplify_omp_ctx *ctx,
-		    enum omp_region_type region_type, enum tree_code code,
-		    tree decl, tree *pd, bool component_ref_p,
-		    unsigned int *flags, tree c,
+static tree *
+build_struct_group (enum omp_region_type region_type, enum tree_code code,
 		    hash_map<tree_operand_hash, tree> *&struct_map_to_clause,
-		    hash_map<tree_operand_hash, tree *> *&struct_seen_clause,
-		    tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p,
-		    bool *cont)
+		    tree *grp_start_p, tree grp_end, tree *inner)
 {
   poly_offset_int coffset;
   poly_int64 cbitpos;
-  tree tree_coffset;
-  tree ocd = OMP_CLAUSE_DECL (c);
+  tree ocd = OMP_CLAUSE_DECL (grp_end);
   bool openmp = !(region_type & ORT_ACC);
+  tree *continue_at = NULL;
 
   while (TREE_CODE (ocd) == ARRAY_REF)
     ocd = TREE_OPERAND (ocd, 0);
@@ -9773,90 +9769,31 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
   if (TREE_CODE (ocd) == INDIRECT_REF)
     ocd = TREE_OPERAND (ocd, 0);
 
-  tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset, &tree_coffset);
-  tree sbase;
+  tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset);
 
-  if (openmp)
-    {
-      if (TREE_CODE (base) == INDIRECT_REF
-	  && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
-	sbase = strip_indirections (base);
-      else
-	sbase = base;
-    }
-  else
-    {
-      sbase = strip_indirections (base);
+  bool ptr = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ALWAYS_POINTER);
+  bool attach_detach = ((OMP_CLAUSE_MAP_KIND (grp_end)
+			 == GOMP_MAP_ATTACH_DETACH)
+			|| (OMP_CLAUSE_MAP_KIND (grp_end)
+			    == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION));
+  bool attach = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ATTACH
+		 || OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_DETACH);
 
-      STRIP_NOPS (sbase);
-    }
-
-  bool do_map_struct = (sbase == decl && !tree_coffset);
-
-  /* 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)
-       : NULL);
-  bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER);
-  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH);
-  bool attach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-		 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH);
-  bool has_attachments = false;
-
-  /* For OpenACC, pointers in structs should trigger an attach action.  */
-  if (attach_detach
-      && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
-	  || code == OMP_TARGET_ENTER_DATA
-	  || code == OMP_TARGET_EXIT_DATA))
-    {
-      /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or
-	 GOMP_MAP_DETACH clause after we have detected a case that needs a
-	 GOMP_MAP_STRUCT mapping added.  */
-      gomp_map_kind k
-	= ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
-	   ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
-      OMP_CLAUSE_SET_MAP_KIND (c, k);
-      has_attachments = true;
-    }
-
-  /* We currently don't handle non-constant offset accesses wrt to
-     GOMP_MAP_STRUCT elements.  */
-  if (!do_map_struct)
-    return NULL_TREE;
-
-  /* Nor for attach_detach for OpenMP.  */
+  /* FIXME: If we're not mapping the base pointer in some other clause on this
+     directive, I think we want to create ALLOC/RELEASE here -- i.e. not
+     early-exit.  */
   if (openmp && attach_detach)
-    {
-      if (DECL_P (decl))
-	{
-	  if (struct_seen_clause == NULL)
-	    struct_seen_clause = new hash_map<tree_operand_hash, tree *>;
-	  if (!struct_seen_clause->get (decl))
-	    struct_seen_clause->put (decl, list_p);
-	}
+    return NULL;
 
-      return NULL_TREE;
-    }
-
-  if ((DECL_P (decl) && (n == NULL || (n->value & GOVD_MAP) == 0))
-      || (!DECL_P (decl)
-	  && (!struct_map_to_clause
-	      || struct_map_to_clause->get (decl) == NULL)))
+  if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
     {
-      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP);
       gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT;
 
       OMP_CLAUSE_SET_MAP_KIND (l, k);
 
       OMP_CLAUSE_DECL (l) = unshare_expr (base);
-      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;
+
       OMP_CLAUSE_SIZE (l)
 	= (!attach ? size_int (1)
 	   : (DECL_P (OMP_CLAUSE_DECL (l))
@@ -9864,19 +9801,17 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 	      : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))));
       if (struct_map_to_clause == NULL)
 	struct_map_to_clause = new hash_map<tree_operand_hash, tree>;
-      struct_map_to_clause->put (decl, l);
+      struct_map_to_clause->put (base, l);
 
       if (ptr || attach_detach)
 	{
 	  tree extra_node;
 	  tree alloc_node
-	    = build_struct_comp_nodes (code, *prev_list_p, c, &extra_node);
+	    = build_struct_comp_nodes (code, *grp_start_p, grp_end,
+				       &extra_node);
 	  OMP_CLAUSE_CHAIN (l) = alloc_node;
 
-	  tree **sc = (struct_seen_clause
-		       ? struct_seen_clause->get (decl)
-		       : NULL);
-	  tree *insert_node_pos = sc ? *sc : prev_list_p;
+	  tree *insert_node_pos = grp_start_p;
 
 	  if (extra_node)
 	    {
@@ -9887,131 +9822,89 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 	    OMP_CLAUSE_CHAIN (alloc_node) = *insert_node_pos;
 
 	  *insert_node_pos = l;
-	  prev_list_p = NULL;
 	}
       else
-	list_p = insert_node_after (l, list_p);
-
-      bool base_ref
-	= (TREE_CODE (base) == INDIRECT_REF
-	   && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
-		== REFERENCE_TYPE)
-	       || ((TREE_CODE (TREE_OPERAND (base, 0)) == INDIRECT_REF)
-		   && (TREE_CODE (TREE_TYPE (TREE_OPERAND
-					      (TREE_OPERAND (base, 0), 0)))
-		       == REFERENCE_TYPE))));
-      bool base_ind = ((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));
-
-	/* 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_ref && code == OMP_TARGET)
 	{
-	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-				      OMP_CLAUSE_MAP);
-	  enum gomp_map_kind mkind
-	    = GOMP_MAP_FIRSTPRIVATE_REFERENCE;
-	  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;
+	  gcc_assert (*grp_start_p == grp_end);
+	  grp_start_p = insert_node_after (l, grp_start_p);
 	}
-      else if (!openmp
-	       && (base_ind || base_ref)
-	       && (region_type & ORT_TARGET))
+
+      tree noind = strip_indirections (base);
+
+      if (!openmp
+	  && (region_type & ORT_TARGET)
+	  && TREE_CODE (noind) == COMPONENT_REF)
 	{
-	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+	  /* The base for this component access 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,
+	     inner GOMP_MAP_STRUCT mapping.
+
+	     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 (grp_end),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
+	  OMP_CLAUSE_DECL (c2) = unshare_expr (noind);
+	  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+	  *inner = c2;
+	  return NULL;
+	}
+
+      tree sdecl = strip_components_and_deref (base);
+
+      if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET))
+	{
+	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+				      OMP_CLAUSE_MAP);
+	  bool base_ref
+	    = (TREE_CODE (base) == INDIRECT_REF
+	       && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
+		    == REFERENCE_TYPE)
+		   || ((TREE_CODE (TREE_OPERAND (base, 0))
+			== INDIRECT_REF)
+		       && (TREE_CODE (TREE_TYPE (TREE_OPERAND
+						  (TREE_OPERAND (base, 0), 0)))
+			   == REFERENCE_TYPE))));
 	  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;
-      if (has_attachments)
-	*flags |= GOVD_MAP_HAS_ATTACHMENTS;
-
-      /* If this is a *pointer-to-struct expression, make sure a
-	 firstprivate map of the base-pointer exists.  */
-      if (openmp
-	  && component_ref_p
-	  && ((TREE_CODE (decl) == MEM_REF
-	       && integer_zerop (TREE_OPERAND (decl, 1)))
-	      || INDIRECT_REF_P (decl))
-	  && DECL_P (TREE_OPERAND (decl, 0))
-	  && !splay_tree_lookup (ctx->variables,
-				 ((splay_tree_key) TREE_OPERAND (decl, 0))))
-	{
-	  decl = TREE_OPERAND (decl, 0);
-	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-	  enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_POINTER;
-	  OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
-	  OMP_CLAUSE_DECL (c2) = decl;
-	  OMP_CLAUSE_SIZE (c2) = size_zero_node;
-	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
-	  OMP_CLAUSE_CHAIN (c) = c2;
+	  OMP_CLAUSE_DECL (c2) = sdecl;
+	  tree baddr = build_fold_addr_expr (base);
+	  baddr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end),
+				    ptrdiff_type_node, baddr);
+	  /* This isn't going to be good enough when we add support for more
+	     complicated lvalue expressions.  FIXME.  */
+	  if (TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE
+	      && TREE_CODE (TREE_TYPE (TREE_TYPE (sdecl))) == POINTER_TYPE)
+	    sdecl = build_simple_mem_ref (sdecl);
+	  tree decladdr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end),
+					    ptrdiff_type_node, sdecl);
+	  OMP_CLAUSE_SIZE (c2)
+	    = fold_build2_loc (OMP_CLAUSE_LOCATION (grp_end), MINUS_EXPR,
+			       ptrdiff_type_node, baddr, decladdr);
+	  /* Insert after struct node.  */
+	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
+	  OMP_CLAUSE_CHAIN (l) = c2;
 	}
 
-      return decl;
+      return NULL;
     }
   else if (struct_map_to_clause)
     {
-      tree *osc = struct_map_to_clause->get (decl);
+      tree *osc = struct_map_to_clause->get (base);
       tree *sc = NULL, *scp = NULL;
-      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 there.)  */
-      if (*sc != c
+      if (*sc != grp_end
 	  && (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)
+      for (; *sc != grp_end; sc = &OMP_CLAUSE_CHAIN (*sc))
+	if ((ptr || attach_detach) && sc == grp_start_p)
 	  break;
 	else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
 		 && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF
@@ -10022,7 +9915,6 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 	    tree sc_decl = OMP_CLAUSE_DECL (*sc);
 	    poly_offset_int offset;
 	    poly_int64 bitpos;
-	    tree tree_offset;
 
 	    if (TREE_CODE (sc_decl) == ARRAY_REF)
 	      {
@@ -10038,8 +9930,7 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 			 == REFERENCE_TYPE))
 	      sc_decl = TREE_OPERAND (sc_decl, 0);
 
-	    tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset,
-						  &tree_offset);
+	    tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset);
 	    if (!base2 || !operand_equal_p (base2, base, 0))
 	      break;
 	    if (scp)
@@ -10049,7 +9940,7 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 		/* This duplicate checking code is currently only enabled for
 		   OpenACC.  */
 		tree d1 = OMP_CLAUSE_DECL (*sc);
-		tree d2 = OMP_CLAUSE_DECL (c);
+		tree d2 = OMP_CLAUSE_DECL (grp_end);
 		while (TREE_CODE (d1) == ARRAY_REF)
 		  d1 = TREE_OPERAND (d1, 0);
 		while (TREE_CODE (d2) == ARRAY_REF)
@@ -10069,10 +9960,10 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 		    break;
 		if (d1 == d2)
 		  {
-		    error_at (OMP_CLAUSE_LOCATION (c),
+		    error_at (OMP_CLAUSE_LOCATION (grp_end),
 			      "%qE appears more than once in map clauses",
-			      OMP_CLAUSE_DECL (c));
-		    return error_mark_node;
+			      OMP_CLAUSE_DECL (grp_end));
+		    return NULL;
 		  }
 	      }
 	    if (maybe_lt (coffset, offset)
@@ -10092,15 +9983,15 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
       if (ptr || attach_detach)
 	{
 	  tree cl = NULL_TREE, extra_node;
-	  tree alloc_node = build_struct_comp_nodes (code, *prev_list_p, c,
-						     &extra_node);
+	  tree alloc_node = build_struct_comp_nodes (code, *grp_start_p,
+						     grp_end, &extra_node);
 	  tree *tail_chain = NULL;
 
 	  /* Here, we have:
 
-	     c : the currently-processed node.
-	     prev_list_p : pointer to the first node in a pointer mapping group
-			   up to and including C.
+	     grp_end : the last (or only) node in this group.
+	     grp_start_p : pointer to the first node in a pointer mapping group
+			   up to and including GRP_END.
 	     sc : pointer to the chain for the end of the struct component
 		  list.
 	     scp : pointer to the chain for the sorted position at which we
@@ -10111,7 +10002,7 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 			  (the end of the struct component list).
 	     extra_node : a newly-synthesized node for an additional indirect
 			  pointer mapping or a Fortran pointer set, if needed.
-	     cl : first node to prepend before prev_list_p.
+	     cl : first node to prepend before grp_start_p.
 	     tail_chain : pointer to chain of last prepended node.
 
 	     The general idea is we move the nodes for this struct mapping
@@ -10147,32 +10038,180 @@  build_struct_group (struct gimplify_omp_ctx *ctx,
 	      tail_chain = &OMP_CLAUSE_CHAIN (alloc_node);
 	    }
 
-	  tree *continue_at
-	    = cl ? move_concat_nodes_after (cl, tail_chain, prev_list_p, c, sc)
-		 : move_nodes_after (prev_list_p, c, sc);
-
-	  prev_list_p = NULL;
-
-	  if (continue_at)
-	    {
-	      list_p = continue_at;
-	      *cont = true;
-	    }
+	  continue_at
+	    = cl ? move_concat_nodes_after (cl, tail_chain, grp_start_p,
+					    grp_end, sc)
+		 : move_nodes_after (grp_start_p, grp_end, sc);
 	}
-      else if (*sc != c)
+      else if (*sc != grp_end)
 	{
-	  if (openmp
-	      && (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
-		  == GS_ERROR))
-	    return error_mark_node;
+	  gcc_assert (*grp_start_p == grp_end);
+
+	  /* We are moving the current node back to a previous struct node:
+	     the node that used to point to the current node will now point to
+	     the next node.  */
+	  continue_at = grp_start_p;
 	  /* 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.  */
-	  move_node_after (c, list_p, sc);
-	  *cont = true;
+	  move_node_after (grp_end, grp_start_p, sc);
 	}
     }
-  return NULL_TREE;
+  return continue_at;
+}
+
+/* Scan through GROUPS, and create sorted structure sibling lists without
+   gimplifying.  */
+
+static bool
+omp_build_struct_sibling_lists (enum tree_code code,
+				enum omp_region_type region_type,
+				vec<omp_mapping_group> *groups,
+				hash_map<tree_operand_hash, omp_mapping_group *>
+				  *grpmap)
+{
+  unsigned i;
+  omp_mapping_group *grp;
+  hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
+  bool success = true;
+  tree *new_next = NULL;
+  tree *tail = &OMP_CLAUSE_CHAIN ((*groups)[groups->length () - 1].grp_end);
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      tree c = grp->grp_end;
+      tree decl = OMP_CLAUSE_DECL (c);
+      tree *grp_start_p = new_next ? new_next : grp->grp_start;
+      tree grp_end = grp->grp_end;
+
+      new_next = NULL;
+
+      if (DECL_P (decl))
+	continue;
+
+      if (OMP_CLAUSE_CHAIN (*grp_start_p)
+	  && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+	{
+	  /* Don't process an array descriptor that isn't inside a derived type
+	     as a struct (the GOMP_MAP_POINTER following will have the form
+	     "var.data", but such mappings are handled specially).  */
+	  tree grpmid = OMP_CLAUSE_CHAIN (*grp_start_p);
+	  if (OMP_CLAUSE_CODE (grpmid) == OMP_CLAUSE_MAP
+	      && OMP_CLAUSE_MAP_KIND (grpmid) == GOMP_MAP_TO_PSET
+	      && DECL_P (OMP_CLAUSE_DECL (grpmid)))
+	    continue;
+	}
+
+      tree d = decl;
+      if (TREE_CODE (d) == ARRAY_REF)
+	{
+	  while (TREE_CODE (d) == ARRAY_REF)
+	    d = TREE_OPERAND (d, 0);
+	  if (TREE_CODE (d) == COMPONENT_REF
+	      && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+	    decl = d;
+	}
+      if (d == decl
+	  && 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))
+	decl = TREE_OPERAND (decl, 0);
+
+      STRIP_NOPS (decl);
+
+      if (TREE_CODE (decl) != COMPONENT_REF)
+	continue;
+
+      omp_mapping_group **wholestruct = NULL;
+      tree wsdecl = omp_containing_struct (OMP_CLAUSE_DECL (c));
+
+      if (!(region_type & ORT_ACC) && wsdecl != OMP_CLAUSE_DECL (c))
+	{
+	  wholestruct = grpmap->get (wsdecl);
+	  if (!wholestruct
+	      && TREE_CODE (wsdecl) == MEM_REF
+	      && integer_zerop (TREE_OPERAND (wsdecl, 1)))
+	    {
+	      tree deref = TREE_OPERAND (wsdecl, 0);
+	      deref = build1 (INDIRECT_REF, TREE_TYPE (wsdecl), deref);
+	      wholestruct = grpmap->get (deref);
+	    }
+	}
+
+      if (wholestruct)
+	{
+
+	  if (*grp_start_p == grp_end)
+	    {
+	      /* Remove the whole of this mapping -- redundant.  */
+	      new_next = grp_start_p;
+	      *grp_start_p = OMP_CLAUSE_CHAIN (grp_end);
+	    }
+
+	  continue;
+	}
+
+      if (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
+	  && code != OACC_UPDATE
+	  && code != OMP_TARGET_UPDATE)
+	{
+	  if (error_operand_p (decl))
+	    {
+	      success = false;
+	      goto error_out;
+	    }
+
+	  tree stype = TREE_TYPE (decl);
+	  if (TREE_CODE (stype) == REFERENCE_TYPE)
+	    stype = TREE_TYPE (stype);
+	  if (TYPE_SIZE_UNIT (stype) == NULL
+	      || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"mapping field %qE of variable length "
+			"structure", OMP_CLAUSE_DECL (c));
+	      success = false;
+	      goto error_out;
+	    }
+
+	  tree inner = NULL_TREE;
+
+	  new_next = build_struct_group (region_type, code,
+					 struct_map_to_clause, grp_start_p,
+					 grp_end, &inner);
+
+	  if (inner)
+	    {
+	      if (new_next && *new_next == NULL_TREE)
+		*new_next = inner;
+	      else
+		*tail = inner;
+
+	      OMP_CLAUSE_CHAIN (inner) = NULL_TREE;
+
+	      omp_mapping_group newgrp;
+	      newgrp.grp_start = new_next ? new_next : tail;
+	      newgrp.grp_end = inner;
+	      newgrp.mark = UNVISITED;
+	      newgrp.sibling = NULL;
+	      newgrp.next = NULL;
+	      groups->safe_push (newgrp);
+
+	      tail = &OMP_CLAUSE_CHAIN (inner);
+	    }
+	}
+    }
+
+error_out:
+  if (struct_map_to_clause)
+    delete struct_map_to_clause;
+
+  return success;
 }
 
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
@@ -10185,9 +10224,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_map<tree_operand_hash, tree *> *struct_seen_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;
@@ -10219,14 +10255,10 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	break;
       }
 
-  /* Topological sorting may fail if we have duplicate nodes, which
-     we should have detected and shown an error for already.  Skip
-     sorting in that case.  */
-  if (!seen_error ()
-      && (code == OMP_TARGET
-	  || code == OMP_TARGET_DATA
-	  || code == OMP_TARGET_ENTER_DATA
-	  || code == OMP_TARGET_EXIT_DATA))
+  if (code == OMP_TARGET
+      || code == OMP_TARGET_DATA
+      || code == OMP_TARGET_ENTER_DATA
+      || code == OMP_TARGET_EXIT_DATA)
     {
       vec<omp_mapping_group> *groups;
       groups = omp_gather_mapping_groups (list_p);
@@ -10234,12 +10266,46 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	{
 	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
 	  grpmap = omp_index_mapping_groups (groups);
-	  omp_mapping_group *outlist
-	    = omp_tsort_mapping_groups (groups, grpmap);
-	  outlist = omp_segregate_mapping_groups (outlist);
-	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+	  omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
+
+	  omp_mapping_group *outlist = NULL;
+
+	  /* Topological sorting may fail if we have duplicate nodes, which
+	     we should have detected and shown an error for already.  Skip
+	     sorting in that case.  */
+	  if (seen_error ())
+	    goto failure;
+
 	  delete grpmap;
 	  delete groups;
+
+	  /* Rebuild now we have struct sibling lists.  */
+	  groups = omp_gather_mapping_groups (list_p);
+	  grpmap = omp_index_mapping_groups (groups);
+
+	  outlist = omp_tsort_mapping_groups (groups, grpmap);
+	  outlist = omp_segregate_mapping_groups (outlist);
+	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+	failure:
+	  delete grpmap;
+	  delete groups;
+	}
+    }
+  else if (region_type & ORT_ACC)
+    {
+      vec<omp_mapping_group> *groups;
+      groups = omp_gather_mapping_groups (list_p);
+      if (groups)
+	{
+	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+	  grpmap = omp_index_mapping_groups (groups);
+
+	  omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
+
+	  delete groups;
+	  delete grpmap;
 	}
     }
 
@@ -10648,6 +10714,28 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
 
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+	    {
+	      tree base = strip_components_and_deref (decl);
+	      if (DECL_P (base))
+		{
+		  decl = base;
+		  splay_tree_node n
+		    = splay_tree_lookup (ctx->variables,
+					 (splay_tree_key) decl);
+		  if (seen_error ()
+		      && n
+		      && (n->value & (GOVD_MAP | GOVD_FIRSTPRIVATE)) != 0)
+		    {
+		      remove = true;
+		      break;
+		    }
+		  flags = GOVD_MAP | GOVD_EXPLICIT;
+
+		  goto do_add_decl;
+		}
+	    }
+
 	  if (TREE_CODE (decl) == TARGET_EXPR)
 	    {
 	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
@@ -10678,143 +10766,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
-	      bool indir_p = false;
-	      bool component_ref_p = false;
-	      tree indir_base = NULL_TREE;
-	      tree orig_decl = decl;
-	      tree decl_ref = NULL_TREE;
-	      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)
-		    {
-		      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;
-			  indir_base = decl;
-			  decl = TREE_OPERAND (decl, 0);
-			  STRIP_NOPS (decl);
-			}
-		      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)
-		    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 = 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))
-		      && (!struct_map_to_clause
-			  || !struct_map_to_clause->get (indir_base)))
-		    {
-		      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;
-		}
 	      /* An "attach/detach" operation on an update directive should
 		 behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that
 		 unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
@@ -10822,91 +10773,49 @@  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 ((((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
-		  && code != OACC_UPDATE
-		  && code != OMP_TARGET_UPDATE)
+
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		{
-		  if (error_operand_p (decl))
+		  if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
+		      == ARRAY_TYPE)
+		    remove = true;
+		  else
 		    {
-		      remove = true;
-		      break;
+		      gomp_map_kind k = ((code == OACC_EXIT_DATA
+					  || code == OMP_TARGET_EXIT_DATA)
+					 ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
+		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		    }
-
-		  tree stype = TREE_TYPE (decl);
-		  if (TREE_CODE (stype) == REFERENCE_TYPE)
-		    stype = TREE_TYPE (stype);
-		  if (TYPE_SIZE_UNIT (stype) == NULL
-		      || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST)
-		    {
-		      error_at (OMP_CLAUSE_LOCATION (c),
-				"mapping field %qE of variable length "
-				"structure", OMP_CLAUSE_DECL (c));
-		      remove = true;
-		      break;
-		    }
-
-		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
-		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
-		    {
-		      /* Error recovery.  */
-		      if (prev_list_p == NULL)
-			{
-			  remove = true;
-			  break;
-			}
-
-		      /* The below prev_list_p based error recovery code is
-			 currently no longer valid for OpenMP.  */
-		      if (code != OMP_TARGET
-			  && code != OMP_TARGET_DATA
-			  && code != OMP_TARGET_UPDATE
-			  && code != OMP_TARGET_ENTER_DATA
-			  && code != OMP_TARGET_EXIT_DATA
-			  && OMP_CLAUSE_CHAIN (*prev_list_p) != c)
-			{
-			  tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
-			  if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
-			    {
-			      remove = true;
-			      break;
-			    }
-			}
-		    }
-		  bool cont = false;
-		  tree add_decl
-		    = build_struct_group (ctx, region_type, code, decl, pd,
-					  component_ref_p, &flags, c,
-					  struct_map_to_clause,
-					  struct_seen_clause, prev_list_p,
-					  list_p, pre_p, &cont);
-		  if (add_decl == error_mark_node)
-		    {
-		      remove = true;
-		      break;
-		    }
-		  else if (add_decl && DECL_P (add_decl))
-		    {
-		      decl = add_decl;
-		      goto do_add_decl;
-		    }
-		  if (cont)
-		    continue;
 		}
-	      else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+
+	      tree cref = decl;
+
+	      while (TREE_CODE (cref) == ARRAY_REF)
+		cref = TREE_OPERAND (cref, 0);
+
+	      if (TREE_CODE (cref) == INDIRECT_REF)
+		cref = TREE_OPERAND (cref, 0);
+
+	      if (TREE_CODE (cref) == COMPONENT_REF)
 		{
-		  gomp_map_kind k = ((code == OACC_EXIT_DATA
-				      || code == OMP_TARGET_EXIT_DATA)
-				     ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
-		  OMP_CLAUSE_SET_MAP_KIND (c, k);
+		  tree base = cref;
+		  while (base && !DECL_P (base))
+		    {
+		      tree innerbase = omp_get_base_pointer (base);
+		      if (!innerbase)
+			break;
+		      base = innerbase;
+		    }
+		  if (base
+		      && DECL_P (base)
+		      && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+		      && POINTER_TYPE_P (TREE_TYPE (base)))
+		    {
+		      splay_tree_node n
+			= splay_tree_lookup (ctx->variables,
+					     (splay_tree_key) base);
+		      n->value |= GOVD_SEEN;
+		    }
 		}
 
 	      if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -11024,24 +10933,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  break;
 		}
 
-	      /* If this was of the form map(*pointer_to_struct), then the
-		 'pointer_to_struct' DECL should be considered deref'ed.  */
-	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
-		   || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c))
-		   || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c)))
-		  && INDIRECT_REF_P (orig_decl)
-		  && DECL_P (TREE_OPERAND (orig_decl, 0))
-		  && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE)
-		{
-		  tree ptr = TREE_OPERAND (orig_decl, 0);
-		  if (!struct_deref_set || !struct_deref_set->contains (ptr))
-		    {
-		      if (!struct_deref_set)
-			struct_deref_set = new hash_set<tree> ();
-		      struct_deref_set->add (ptr);
-		    }
-		}
-
 	      if (!remove
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -11058,28 +10949,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	      break;
 	    }
-	  else
-	    {
-	      /* DECL_P (decl) == true  */
-	      tree *sc;
-	      if (struct_map_to_clause
-		  && (sc = struct_map_to_clause->get (decl)) != NULL
-		  && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
-		  && decl == OMP_CLAUSE_DECL (*sc))
-		{
-		  /* We have found a map of the whole structure after a
-		     leading GOMP_MAP_STRUCT has been created, so refill the
-		     leading clause into a map of the whole structure
-		     variable, and remove the current one.
-		     TODO: we should be able to remove some maps of the
-		     following structure element maps if they are of
-		     compatible TO/FROM/ALLOC type.  */
-		  OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
-		  OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
-		  remove = true;
-		  break;
-		}
-	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -11721,12 +11590,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
   ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
-  if (struct_seen_clause)
-    delete struct_seen_clause;
-  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
@@ -11875,8 +11738,6 @@  gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     return 0;
   if ((flags & GOVD_SEEN) == 0)
     return 0;
-  if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
-    return 0;
   if (flags & GOVD_DEBUG_PRIVATE)
     {
       gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
index e0c11570f5d..9993768ef20 100644
--- a/gcc/testsuite/g++.dg/goacc/member-array-acc.C
+++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
@@ -10,4 +10,4 @@  struct Foo {
 };
 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" } } */
+/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:this->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach: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..432f02614d8 100644
--- a/gcc/testsuite/g++.dg/gomp/target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -33,4 +33,6 @@  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 S \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 150f286e312..bff7fa7c669 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,7 +87,7 @@  int main (void)
   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\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* { 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\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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" } } */
 
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index 679c85a54dd..cc08e7e8693 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\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) 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:v \[len: [0-9]+\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 8166f43ad42..9ade3cc0b2b 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@  int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) 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]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */