[08/14] OpenMP: Pointers and member mappings

diff --git a/gcc/fortran/dependency.cc b/gcc/fortran/dependency.cc
index 9117825ee6e..5c31044b4d3 100644
--- a/gcc/fortran/dependency.cc
+++ b/gcc/fortran/dependency.cc
@@ -2339,3 +2339,131 @@  gfc_dep_resolver (gfc_ref *lref, gfc_ref *rref, gfc_reverse *reverse,
 
   return fin_dep == GFC_DEP_OVERLAP;
 }
+
+/* Check if two refs are equal, for the purposes of checking if one might be
+   the base of the other for OpenMP (target directives).  Derived from
+   gfc_dep_resolver.  This function is stricter, e.g. indices arr(i) and
+   arr(j) compare as non-equal.  */
+
+bool
+gfc_omp_expr_prefix_same (gfc_expr *lexpr, gfc_expr *rexpr)
+{
+  gfc_ref *lref, *rref;
+
+  if (lexpr->symtree && rexpr->symtree)
+    {
+      /* See are_identical_variables above.  */
+      if (lexpr->symtree->n.sym->attr.dummy
+	  && rexpr->symtree->n.sym->attr.dummy)
+	{
+	  /* Dummy arguments: Only check for equal names.  */
+	  if (lexpr->symtree->n.sym->name != rexpr->symtree->n.sym->name)
+	    return false;
+	}
+      else
+	{
+	  if (lexpr->symtree->n.sym != rexpr->symtree->n.sym)
+	    return false;
+	}
+    }
+  else if (lexpr->base_expr && rexpr->base_expr)
+    {
+      if (gfc_dep_compare_expr (lexpr->base_expr, rexpr->base_expr) != 0)
+	return false;
+    }
+  else
+    return false;
+
+  lref = lexpr->ref;
+  rref = rexpr->ref;
+
+  while (lref && rref)
+    {
+      gfc_dependency fin_dep = GFC_DEP_EQUAL;
+
+      if (lref && lref->type == REF_COMPONENT && lref->u.c.component
+	  && strcmp (lref->u.c.component->name, "_data") == 0)
+	lref = lref->next;
+
+      if (rref && rref->type == REF_COMPONENT && rref->u.c.component
+	  && strcmp (rref->u.c.component->name, "_data") == 0)
+	rref = rref->next;
+
+      gcc_assert (lref->type == rref->type);
+
+      switch (lref->type)
+	{
+	case REF_COMPONENT:
+	  if (lref->u.c.component != rref->u.c.component)
+	    return false;
+	  break;
+
+	case REF_ARRAY:
+	  if (ref_same_as_full_array (lref, rref))
+	    break;
+	  if (ref_same_as_full_array (rref, lref))
+	    break;
+
+	  if (lref->u.ar.dimen != rref->u.ar.dimen)
+	    {
+	      if (lref->u.ar.type == AR_FULL
+		  && gfc_full_array_ref_p (rref, NULL))
+		break;
+	      if (rref->u.ar.type == AR_FULL
+		  && gfc_full_array_ref_p (lref, NULL))
+		break;
+	      return false;
+	    }
+
+	  for (int n = 0; n < lref->u.ar.dimen; n++)
+	    {
+	      if (lref->u.ar.dimen_type[n] == DIMEN_VECTOR
+		  && rref->u.ar.dimen_type[n] == DIMEN_VECTOR
+		  && gfc_dep_compare_expr (lref->u.ar.start[n],
+					   rref->u.ar.start[n]) == 0)
+		continue;
+	      if (lref->u.ar.dimen_type[n] == DIMEN_RANGE
+		  && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+		fin_dep = check_section_vs_section (&lref->u.ar, &rref->u.ar,
+						    n);
+	      else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+		       && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+		fin_dep = gfc_check_element_vs_section (lref, rref, n);
+	      else if (rref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+		       && lref->u.ar.dimen_type[n] == DIMEN_RANGE)
+		fin_dep = gfc_check_element_vs_section (rref, lref, n);
+	      else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+		       && rref->u.ar.dimen_type[n] == DIMEN_ELEMENT)
+		{
+		  gfc_array_ref l_ar = lref->u.ar;
+		  gfc_array_ref r_ar = rref->u.ar;
+		  gfc_expr *l_start = l_ar.start[n];
+		  gfc_expr *r_start = r_ar.start[n];
+		  int i = gfc_dep_compare_expr (r_start, l_start);
+		  if (i == 0)
+		    fin_dep = GFC_DEP_EQUAL;
+		  else
+		    return false;
+		}
+	      else
+		return false;
+	      if (n + 1 < lref->u.ar.dimen
+		  && fin_dep != GFC_DEP_EQUAL)
+		return false;
+	    }
+
+	  if (fin_dep != GFC_DEP_EQUAL
+	      && fin_dep != GFC_DEP_OVERLAP)
+	    return false;
+
+	  break;
+
+	default:
+	  gcc_unreachable ();
+	}
+      lref = lref->next;
+      rref = rref->next;
+    }
+
+  return true;
+}
diff --git a/gcc/fortran/dependency.h b/gcc/fortran/dependency.h
index b1150316f5d..d5dffdd0bcb 100644
--- a/gcc/fortran/dependency.h
+++ b/gcc/fortran/dependency.h
@@ -40,5 +40,6 @@  int gfc_expr_is_one (gfc_expr *, int);
 int gfc_dep_resolver (gfc_ref *, gfc_ref *, gfc_reverse *,
 		      bool identical = false);
 int gfc_are_equivalenced_arrays (gfc_expr *, gfc_expr *);
+bool gfc_omp_expr_prefix_same (gfc_expr *, gfc_expr *);
 
 gfc_expr * gfc_discard_nops (gfc_expr *);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 9c8d88669a0..23afacd015e 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1378,6 +1378,7 @@  typedef struct gfc_omp_namelist
     {
       struct gfc_omp_namelist_udr *udr;
       gfc_namespace *ns;
+      struct gfc_omp_namelist *duplicate_of;
     } u2;
   struct gfc_symbol *memspace_sym;
   struct gfc_symbol *traits_sym;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a055aa5a61a..a108f718ffa 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -49,6 +49,7 @@  along with GCC; see the file COPYING3.  If not see
 #include "stor-layout.h"
 #include "gimple-iterator.h"
 #include "gimplify-me.h"
+#include "dependency.h"
 
 #undef GCC_DIAG_STYLE
 #define GCC_DIAG_STYLE __gcc_tdiag__
@@ -3939,36 +3940,23 @@  gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
     }
   if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
     {
-      tree desc_node;
       tree type = TREE_TYPE (decl);
       ptr2 = gfc_conv_descriptor_data_get (decl);
-      desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
-      OMP_CLAUSE_DECL (desc_node) = decl;
-      OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
-      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
+      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_DECL (node2) = decl;
+      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
+	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
+	  || op == EXEC_OMP_TARGET_EXIT_DATA)
 	{
-	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_DELETE);
-	  node2 = desc_node;
-	}
-      else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
-	       || op == EXEC_OMP_TARGET_EXIT_DATA)
-	{
-	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_RELEASE);
-	  node2 = desc_node;
-	}
-      else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
-	{
-	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
-	  node2 = node;
-	  node = desc_node;  /* Needs to come first.  */
+	  gomp_map_kind map_kind
+	    = (op == EXEC_OMP_TARGET_EXIT_DATA) ? GOMP_MAP_RELEASE
+						: OMP_CLAUSE_MAP_KIND (node);
+	  OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+	  OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;
 	}
       else
-	{
-	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
-	  node2 = desc_node;
-	}
-      if (op == EXEC_OMP_TARGET_EXIT_DATA)
-	return;
+	OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
       node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
       OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
@@ -4072,6 +4060,73 @@  handle_iterator (gfc_namespace *ns, stmtblock_t *iter_block, tree block)
   return list;
 }
 
+/* To alleviate quadratic behaviour in checking each entry of a
+   gfc_omp_namelist against every other entry, we build a hashtable indexed by
+   gfc_symbol pointer, which we can use in the usual case that a map
+   expression has a symbol as its root term.  Return a namelist based on the
+   root symbol used by N, building a new table in SYM_ROOTED_NL using the
+   gfc_omp_namelist N2 (all clauses) if we haven't done so already.  */
+
+static gfc_omp_namelist *
+get_symbol_rooted_namelist (hash_map<gfc_symbol *,
+				     gfc_omp_namelist *> *&sym_rooted_nl,
+			    gfc_omp_namelist *n,
+			    gfc_omp_namelist *n2, bool *sym_based)
+{
+  /* Early-out if we have a NULL clause list (e.g. for OpenACC).  */
+  if (!n2)
+    return NULL;
+
+  gfc_symbol *use_sym = NULL;
+
+  /* We're only interested in cases where we have an expression, e.g. a
+     component access.  */
+  if (n->expr && n->expr->expr_type == EXPR_VARIABLE && n->expr->symtree)
+    use_sym = n->expr->symtree->n.sym;
+
+  *sym_based = false;
+
+  if (!use_sym)
+    return n2;
+
+  if (!sym_rooted_nl)
+    {
+      sym_rooted_nl = new hash_map<gfc_symbol *, gfc_omp_namelist *> ();
+
+      for (; n2 != NULL; n2 = n2->next)
+	{
+	  if (!n2->expr
+	      || n2->expr->expr_type != EXPR_VARIABLE
+	      || !n2->expr->symtree)
+	    continue;
+
+	  gfc_omp_namelist *nl_copy = gfc_get_omp_namelist ();
+	  memcpy (nl_copy, n2, sizeof *nl_copy);
+	  nl_copy->u2.duplicate_of = n2;
+	  nl_copy->next = NULL;
+
+	  gfc_symbol *idx_sym = n2->expr->symtree->n.sym;
+
+	  bool existed;
+	  gfc_omp_namelist *&entry
+	    = sym_rooted_nl->get_or_insert (idx_sym, &existed);
+	  if (existed)
+	    nl_copy->next = entry;
+	  entry = nl_copy;
+	}
+    }
+
+  gfc_omp_namelist **n2_sym = sym_rooted_nl->get (use_sym);
+
+  if (n2_sym)
+    {
+      *sym_based = true;
+      return *n2_sym;
+    }
+
+  return NULL;
+}
+
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		       locus where, bool declare_simd = false,
@@ -4089,6 +4144,8 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
   if (clauses == NULL)
     return NULL_TREE;
 
+  hash_map<gfc_symbol *, gfc_omp_namelist *> *sym_rooted_nl = NULL;
+
   for (list = 0; list < OMP_LIST_NUM; list++)
     {
       gfc_omp_namelist *n = clauses->lists[list];
@@ -5202,6 +5259,54 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      goto finalize_map_clause;
 			    }
 
+			  gfc_omp_namelist *n2
+			    = openacc ? NULL : clauses->lists[OMP_LIST_MAP];
+
+			  bool sym_based;
+			  n2 = get_symbol_rooted_namelist (sym_rooted_nl, n,
+							   n2, &sym_based);
+
+			  /* If the last reference is a pointer to a derived
+			     type ("foo%dt_ptr"), check if any subcomponents
+			     of the same derived type member are being mapped
+			     elsewhere in the clause list ("foo%dt_ptr%x",
+			     etc.).  If we have such subcomponent mappings,
+			     we only create an ALLOC node for the pointer
+			     itself, and inhibit mapping the whole derived
+			     type.  */
+
+			  for (; n2 != NULL; n2 = n2->next)
+			    {
+			      if ((!sym_based && n == n2)
+				  || (sym_based && n == n2->u2.duplicate_of)
+				  || !n2->expr)
+				continue;
+
+			      if (!gfc_omp_expr_prefix_same (n->expr,
+							     n2->expr))
+				continue;
+
+			      gfc_ref *ref1 = n->expr->ref;
+			      gfc_ref *ref2 = n2->expr->ref;
+
+			      while (ref1->next && ref2->next)
+				{
+				  ref1 = ref1->next;
+				  ref2 = ref2->next;
+				}
+
+			      if (ref2->next)
+				{
+				  inner = build_fold_addr_expr (inner);
+				  OMP_CLAUSE_SET_MAP_KIND (node,
+							   GOMP_MAP_ALLOC);
+				  OMP_CLAUSE_DECL (node) = inner;
+				  OMP_CLAUSE_SIZE (node)
+				    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+				  goto finalize_map_clause;
+				}
+			    }
+
 			  tree data, size;
 
 			  if (lastref->u.c.component->ts.type == BT_CLASS)
@@ -5273,7 +5378,6 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
 			{
 			  gomp_map_kind map_kind;
-			  tree desc_node;
 			  tree type = TREE_TYPE (inner);
 			  tree ptr = gfc_conv_descriptor_data_get (inner);
 			  ptr = build_fold_indirect_ref (ptr);
@@ -5316,21 +5420,19 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  OMP_CLAUSE_SIZE (node)
 			    = fold_build2 (MULT_EXPR, gfc_array_index_type,
 					   OMP_CLAUSE_SIZE (node), elemsz);
-			  desc_node = build_omp_clause (input_location,
-							OMP_CLAUSE_MAP);
-			  if (openacc)
-			    OMP_CLAUSE_SET_MAP_KIND (desc_node,
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  if (openacc
+			      || (map_kind != GOMP_MAP_RELEASE
+				  && map_kind != GOMP_MAP_DELETE))
+			    OMP_CLAUSE_SET_MAP_KIND (node2,
 						     GOMP_MAP_TO_PSET);
 			  else
-			    OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
-			  OMP_CLAUSE_DECL (desc_node) = inner;
-			  OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
-			  if (openacc)
-			    node2 = desc_node;
-			  else
+			    OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+			  OMP_CLAUSE_DECL (node2) = inner;
+			  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+			  if (!openacc)
 			    {
-			      node2 = node;
-			      node = desc_node;  /* Put first.  */
 			      if (n->expr->ts.type == BT_DERIVED
 				  && n->expr->ts.u.derived->attr.alloc_comp)
 				{
@@ -5338,18 +5440,73 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 				     in gfc_omp_deep_mapping{,_p,_cnt}; force
 				     evaluate to ensure that it is
 				     not gimplified + is a decl.  */
-				  tree tmp = OMP_CLAUSE_SIZE (node2);
+				  tree tmp = OMP_CLAUSE_SIZE (node);
 				  tree var = gfc_create_var (TREE_TYPE (tmp),
 							     NULL);
 				  gfc_add_modify_loc (input_location, block,
 						      var, tmp);
-				  OMP_CLAUSE_SIZE (node2) = var;
+				  OMP_CLAUSE_SIZE (node) = var;
 				  gfc_allocate_lang_decl (var);
 				  GFC_DECL_SAVED_DESCRIPTOR (var) = inner;
 				}
+
+			      gfc_omp_namelist *n2
+				= clauses->lists[OMP_LIST_MAP];
+
+			      /* If we don't have a mapping of a smaller part
+				 of the array -- or we can't prove that we do
+				 statically -- set this flag.  If there is a
+				 mapping of a smaller part of the array after
+				 all, this will turn into a no-op at
+				 runtime.  */
+			      OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1;
+
+			      bool sym_based;
+			      n2 = get_symbol_rooted_namelist (sym_rooted_nl,
+							       n, n2,
+							       &sym_based);
+
+			      bool drop_mapping = false;
+
+			      for (; n2 != NULL; n2 = n2->next)
+				{
+				  if ((!sym_based && n == n2)
+				      || (sym_based && n == n2->u2.duplicate_of)
+				      || !n2->expr)
+				    continue;
+
+				  if (!gfc_omp_expr_prefix_same (n->expr,
+								 n2->expr))
+				    continue;
+
+				  gfc_ref *ref1 = n->expr->ref;
+				  gfc_ref *ref2 = n2->expr->ref;
+
+				  /* We know ref1 and ref2 overlap.  We're
+				     interested in whether ref2 describes a
+				     smaller part of the array than ref1, which
+				     we already know refers to the full
+				     array.  */
+
+				  while (ref1->next && ref2->next)
+				    {
+				      ref1 = ref1->next;
+				      ref2 = ref2->next;
+				    }
+
+				  if (ref2->next
+				      || (ref2->type == REF_ARRAY
+					  && (ref2->u.ar.type == AR_ELEMENT
+					      || (ref2->u.ar.type
+						  == AR_SECTION))))
+				    {
+				      drop_mapping = true;
+				      break;
+				    }
+				}
+			      if (drop_mapping)
+				continue;
 			    }
-			  if (op == EXEC_OMP_TARGET_EXIT_DATA)
-			    goto finalize_map_clause;
 			  node3 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
 			  OMP_CLAUSE_SET_MAP_KIND (node3,
@@ -5527,6 +5684,23 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	}
     }
 
+  /* Free hashmap if we built it.  */
+  if (sym_rooted_nl)
+    {
+      typedef hash_map<gfc_symbol *, gfc_omp_namelist *>::iterator hti;
+      for (hti it = sym_rooted_nl->begin (); it != sym_rooted_nl->end (); ++it)
+	{
+	  gfc_omp_namelist *&nl = (*it).second;
+	  while (nl)
+	    {
+	      gfc_omp_namelist *next = nl->next;
+	      free (nl);
+	      nl = next;
+	    }
+	}
+      delete sym_rooted_nl;
+    }
+
   if (clauses->if_expr)
     {
       tree if_var;
@@ -6397,7 +6571,7 @@  gfc_trans_oacc_executable_directive (gfc_code *code)
 
   gfc_start_block (&block);
   oacc_clauses = gfc_trans_omp_clauses (&block, clauses, code->loc,
-					false, true);
+					false, true, code->op);
   stmt = build1_loc (input_location, construct_code, void_type_node, 
 		     oacc_clauses);
   gfc_add_expr_to_block (&block, stmt);
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 6a43c792450..da81582da1c 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8865,6 +8865,25 @@  gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
   return 1;
 }
 
+/* True if mapping node C maps, or unmaps, a (Fortran) array descriptor.  */
+
+static bool
+omp_map_clause_descriptor_p (tree c)
+{
+  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+    return false;
+
+  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
+    return true;
+
+  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
+       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE)
+      && OMP_CLAUSE_RELEASE_DESCRIPTOR (c))
+    return true;
+
+  return false;
+}
+
 /* For a set of mappings describing an array section pointed to by a struct
    (or derived type, etc.) component, create an "alloc" or "release" node to
    insert into a list following a GOMP_MAP_STRUCT node.  For some types of
@@ -8900,9 +8919,7 @@  build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
   if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
     grp_mid = OMP_CLAUSE_CHAIN (grp_start);
 
-  if (grp_mid
-      && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
-      && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+  if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
     OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (grp_mid);
   else
     OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
@@ -9088,7 +9105,7 @@  omp_get_attachment (omp_mapping_group *grp)
 	return NULL_TREE;
 
       node = OMP_CLAUSE_CHAIN (node);
-      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+      if (node && omp_map_clause_descriptor_p (node))
 	{
 	  gcc_assert (node != grp->grp_end);
 	  node = OMP_CLAUSE_CHAIN (node);
@@ -9181,7 +9198,7 @@  omp_group_last (tree *start_p)
 		 || (OMP_CLAUSE_MAP_KIND (nc)
 		     == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
-		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+		 || omp_map_clause_descriptor_p (nc)))
 	{
 	  grp_last_p = &OMP_CLAUSE_CHAIN (c);
 	  c = nc;
@@ -9337,32 +9354,31 @@  omp_group_base (omp_mapping_group *grp, unsigned int *chained,
 	return node;
 
       node = OMP_CLAUSE_CHAIN (node);
-      if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+      if (!node)
+	internal_error ("unexpected mapping node");
+      if (omp_map_clause_descriptor_p (node))
 	{
 	  if (node == grp->grp_end)
 	    return *grp->grp_start;
 	  node = OMP_CLAUSE_CHAIN (node);
 	}
-      if (node)
-	switch (OMP_CLAUSE_MAP_KIND (node))
-	  {
-	  case GOMP_MAP_POINTER:
-	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
-	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
-	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
-	    *firstprivate = OMP_CLAUSE_DECL (node);
-	    return *grp->grp_start;
+      switch (OMP_CLAUSE_MAP_KIND (node))
+	{
+	case GOMP_MAP_POINTER:
+	case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+	case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	  *firstprivate = OMP_CLAUSE_DECL (node);
+	  return *grp->grp_start;
 
-	  case GOMP_MAP_ALWAYS_POINTER:
-	  case GOMP_MAP_ATTACH_DETACH:
-	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
-	    return *grp->grp_start;
+	case GOMP_MAP_ALWAYS_POINTER:
+	case GOMP_MAP_ATTACH_DETACH:
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  return *grp->grp_start;
 
-	  default:
-	    internal_error ("unexpected mapping node");
-	  }
-      else
-	internal_error ("unexpected mapping node");
+	default:
+	  internal_error ("unexpected mapping node");
+	}
       return error_mark_node;
 
     case GOMP_MAP_TO_PSET:
@@ -9712,18 +9728,45 @@  omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
 static omp_mapping_group *
 omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
 			  hash_map<tree_operand_hash_no_se, omp_mapping_group *>
-			    *grpmap)
+			    *grpmap,
+			  bool enter_exit_data)
 {
   omp_mapping_group *grp, *outlist = NULL, **cursor;
   unsigned int i;
+  bool saw_runtime_implicit = false;
 
   cursor = &outlist;
 
   FOR_EACH_VEC_ELT (*groups, i, grp)
     {
       if (grp->mark != PERMANENT)
-	if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
-	  return NULL;
+	{
+	  if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+	    {
+	      saw_runtime_implicit = true;
+	      continue;
+	    }
+	  if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+	    return NULL;
+	}
+    }
+
+  if (!saw_runtime_implicit)
+    return outlist;
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      if (grp->mark != PERMANENT
+	  && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+	{
+	  /* Clear the flag for enter/exit data because it is currently
+	     meaningless for those operations in libgomp.  */
+	  if (enter_exit_data)
+	    OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start) = 0;
+
+	  if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+	    return NULL;
+	}
     }
 
   return outlist;
@@ -10126,6 +10169,11 @@  omp_check_mapping_compatibility (location_t loc,
    mapping.  However, if we have a reference to pointer, make other appropriate
    adjustments to the mapping nodes instead.
 
+   If we have an ATTACH_DETACH node with a Fortran pointer-set (array
+   descriptor) mapping for a derived-type component, and we're also mapping the
+   whole of the derived-type variable on another clause, the pointer-set
+   mapping is removed.
+
    If we have a component access but we're also mapping the whole of the
    containing struct, drop the former access.
 
@@ -10305,6 +10353,17 @@  omp_resolve_clause_dependencies (enum tree_code code,
 		   GOMP_MAP_ATTACH_ZLAS for it.  */
 		if (!base_mapped_to && referenced_ptr_node)
 		  OMP_CLAUSE_SET_MAP_KIND (referenced_ptr_node, zlas_kind);
+
+		omp_mapping_group *struct_group;
+		tree desc;
+		if ((desc = OMP_CLAUSE_CHAIN (*grp->grp_start))
+		    && omp_map_clause_descriptor_p (desc)
+		    && omp_mapped_by_containing_struct (grpmap, decl,
+							&struct_group))
+		  /* If we have a pointer set but we're mapping (or unmapping)
+		     the whole of the containing struct, we can remove the
+		     pointer set mapping.  */
+		  OMP_CLAUSE_CHAIN (*grp->grp_start) = OMP_CLAUSE_CHAIN (desc);
 	      }
 	    else if (TREE_CODE (TREE_TYPE (base_ptr)) == REFERENCE_TYPE
 		     && (TREE_CODE (TREE_TYPE (TREE_TYPE (base_ptr)))
@@ -10752,11 +10811,17 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
      for the purposes of gathering sibling lists, etc.  */
   /* gcc_assert (base == addr_tokens[base_token]->expr);  */
 
-  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 has_descriptor = false;
+  if (OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+    {
+      tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+      if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
+	has_descriptor = true;
+    }
 
   if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
     {
@@ -10779,7 +10844,18 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
 	 GOMP_MAP_STRUCT into the middle of the old one.  */
       tree *insert_node_pos = reprocessing_struct ? *added_tail : grp_start_p;
 
-      if (ptr || attach_detach)
+      if (has_descriptor)
+	{
+	  tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+	  if (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+	    OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+	  tree sc = *insert_node_pos;
+	  OMP_CLAUSE_CHAIN (l) = desc;
+	  OMP_CLAUSE_CHAIN (*grp_start_p) = OMP_CLAUSE_CHAIN (desc);
+	  OMP_CLAUSE_CHAIN (desc) = sc;
+	  *insert_node_pos = l;
+	}
+      else if (attach_detach)
 	{
 	  tree extra_node;
 	  tree alloc_node
@@ -11010,7 +11086,7 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
 	  || OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_ATTACH_DETACH)
 	sc = &OMP_CLAUSE_CHAIN (*sc);
       for (i = 0; i < elems; i++, sc = &OMP_CLAUSE_CHAIN (*sc))
-	if ((ptr || attach_detach) && sc == grp_start_p)
+	if (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
@@ -11066,7 +11142,7 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
 		|| (known_eq (coffset, offset)
 		    && maybe_lt (cbitpos, bitpos)))
 	      {
-		if (ptr || attach_detach)
+		if (attach_detach)
 		  scp = sc;
 		else
 		  break;
@@ -11082,7 +11158,9 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
 	     the list manipulation below.  We only need to handle the (pointer
 	     or reference) attach/detach case.  */
 	  tree extra_node, alloc_node;
-	  if (attach_detach)
+	  if (has_descriptor)
+	    gcc_unreachable ();
+	  else if (attach_detach)
 	    alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
 						      grp_end, &extra_node);
 	  else
@@ -11115,7 +11193,17 @@  omp_accumulate_sibling_list (enum omp_region_type region_type,
 	  return NULL;
 	}
 
-      if (ptr || attach_detach)
+      if (has_descriptor)
+	{
+	  tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+	  if (code == OMP_TARGET_EXIT_DATA
+	      || code == OACC_EXIT_DATA)
+	    OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+	  omp_siblist_move_node_after (desc,
+				       &OMP_CLAUSE_CHAIN (*grp_start_p),
+				       scp ? scp : sc);
+	}
+      else if (attach_detach)
 	{
 	  tree cl = NULL_TREE, extra_node;
 	  tree alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
@@ -11260,8 +11348,7 @@  omp_build_struct_sibling_lists (enum tree_code code,
 	     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
+	  if (omp_map_clause_descriptor_p (grpmid)
 	      && DECL_P (OMP_CLAUSE_DECL (grpmid)))
 	    continue;
 	}
@@ -11537,6 +11624,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 					  list_p);
 
 	  omp_mapping_group *outlist = NULL;
+	  bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+			     || code == OMP_TARGET_EXIT_DATA);
 
 	  /* Topological sorting may fail if we have duplicate nodes, which
 	     we should have detected and shown an error for already.  Skip
@@ -11551,7 +11640,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  groups = omp_gather_mapping_groups (list_p);
 	  grpmap = omp_index_mapping_groups (groups);
 
-	  outlist = omp_tsort_mapping_groups (groups, grpmap);
+	  outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit);
 	  outlist = omp_segregate_mapping_groups (outlist);
 	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
 
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 b/gcc/testsuite/gfortran.dg/gomp/map-9.f90
index f930a49d9ff..8c8d4f7c57c 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-9.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-9.f90
@@ -2,7 +2,7 @@ 
 
 ! PR fortran/108545
 
-! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x\.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(attach:x\.a\.data \\\[bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x\.a \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(attach:x\.a\.data \\\[bias: 0\\\]\\)" "omplower" } }
 
 program p
    type t
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90
new file mode 100644
index 00000000000..5aaaf53e697
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90
@@ -0,0 +1,57 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+integer, pointer :: arr3(:)
+integer, pointer :: arr4(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+tv%arr3 => arr
+tv%arr4 => arr
+
+!$omp target enter data map(to: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(release:tv\.arr1 \[len: [0-9]+\]\) map\(detach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr2) map(to: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr2) map(from: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr2 \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(to:tv \[len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(from:tv \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)[_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr4 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr4 \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90
new file mode 100644
index 00000000000..197888a4336
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90
@@ -0,0 +1,40 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+
+!$omp target map(tv%arr1)
+tv%arr1(1) = tv%arr1(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\[implicit\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target map(tv%arr2) map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+!$omp target map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} 2 "gimple" } }
+
+!$omp target map(tv, tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(tofrom:tv \[len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
diff --git a/gcc/tree.h b/gcc/tree.h
index 58f2e2e20cb..8e2e3654b56 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1797,6 +1797,10 @@  class auto_suppress_location_wrappers
    same directive.  */
 #define OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED(NODE) \
   TREE_STATIC (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if this is an attach/detach that is indirect, i.e. not applied to
+   a structure member.  */
+#define OMP_CLAUSE_RELEASE_DESCRIPTOR(NODE) \
+  TREE_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
 
 /* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP
    lowering.  */
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90
new file mode 100644
index 00000000000..02f08c52a8c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90
@@ -0,0 +1,108 @@ 
+! { dg-do run }
+
+program myprog
+type u
+  integer, dimension (:), pointer :: tarr1
+  integer, dimension (:), pointer :: tarr2
+  integer, dimension (:), pointer :: tarr3
+end type u
+
+type(u) :: myu1, myu2, myu3
+
+integer, dimension (12), target :: myarray1
+integer, dimension (12), target :: myarray2
+integer, dimension (12), target :: myarray3
+integer, dimension (12), target :: myarray4
+integer, dimension (12), target :: myarray5
+integer, dimension (12), target :: myarray6
+integer, dimension (12), target :: myarray7
+integer, dimension (12), target :: myarray8
+integer, dimension (12), target :: myarray9
+
+myu1%tarr1 => myarray1
+myu1%tarr2 => myarray2
+myu1%tarr3 => myarray3
+myu2%tarr1 => myarray4
+myu2%tarr2 => myarray5
+myu2%tarr3 => myarray6
+myu3%tarr1 => myarray7
+myu3%tarr2 => myarray8
+myu3%tarr3 => myarray9
+
+myu1%tarr1 = 0
+myu1%tarr2 = 0
+myu1%tarr3 = 0
+myu2%tarr1 = 0
+myu2%tarr2 = 0
+myu2%tarr3 = 0
+myu3%tarr1 = 0
+myu3%tarr2 = 0
+myu3%tarr3 = 0
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(:)) &
+!$omp&       map(to:myu1%tarr2) map(tofrom:myu1%tarr2(:)) &
+!$omp&       map(to:myu1%tarr3) map(tofrom:myu1%tarr3(:)) &
+!$omp&       map(to:myu2%tarr1) map(tofrom:myu2%tarr1(:)) &
+!$omp&       map(to:myu2%tarr2) map(tofrom:myu2%tarr2(:)) &
+!$omp&       map(to:myu2%tarr3) map(tofrom:myu2%tarr3(:)) &
+!$omp&       map(to:myu3%tarr1) map(tofrom:myu3%tarr1(:)) &
+!$omp&       map(to:myu3%tarr2) map(tofrom:myu3%tarr2(:)) &
+!$omp&       map(to:myu3%tarr3) map(tofrom:myu3%tarr3(:))
+myu1%tarr1(1) = myu1%tarr1(1) + 1
+myu2%tarr1(1) = myu2%tarr1(1) + 1
+myu3%tarr1(1) = myu3%tarr1(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1:2)) &
+!$omp&       map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1:2)) &
+!$omp&       map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1:2)) &
+!$omp&       map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1:2)) &
+!$omp&       map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1:2)) &
+!$omp&       map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1:2)) &
+!$omp&       map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1:2)) &
+!$omp&       map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1:2)) &
+!$omp&       map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1:2))
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1)) &
+!$omp&       map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1)) &
+!$omp&       map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1)) &
+!$omp&       map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1)) &
+!$omp&       map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1)) &
+!$omp&       map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1)) &
+!$omp&       map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1)) &
+!$omp&       map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1)) &
+!$omp&       map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1))
+myu1%tarr3(1) = myu1%tarr3(1) + 1
+myu2%tarr3(1) = myu2%tarr3(1) + 1
+myu3%tarr3(1) = myu3%tarr3(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu1%tarr1) &
+!$omp&       map(tofrom:myu1%tarr2) &
+!$omp&       map(tofrom:myu1%tarr3) &
+!$omp&       map(tofrom:myu2%tarr1) &
+!$omp&       map(tofrom:myu2%tarr2) &
+!$omp&       map(tofrom:myu2%tarr3) &
+!$omp&       map(tofrom:myu3%tarr1) &
+!$omp&       map(tofrom:myu3%tarr2) &
+!$omp&       map(tofrom:myu3%tarr3)
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+if (myu1%tarr1(1).ne.1) stop 1
+if (myu2%tarr1(1).ne.1) stop 2
+if (myu3%tarr1(1).ne.1) stop 3
+if (myu1%tarr2(1).ne.2) stop 4
+if (myu2%tarr2(1).ne.2) stop 5
+if (myu3%tarr2(1).ne.2) stop 6
+if (myu1%tarr3(1).ne.1) stop 7
+if (myu2%tarr3(1).ne.1) stop 8
+if (myu3%tarr3(1).ne.1) stop 9
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90
new file mode 100644
index 00000000000..318e77ea44f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90
@@ -0,0 +1,62 @@ 
+! { dg-do run }
+
+module mymod
+type G
+integer :: x, y
+integer, pointer :: arr(:)
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+integer, target :: arr1(10)
+integer, target :: arr2(10)
+integer, target :: arr3(10)
+type(G), dimension(3) :: gvar
+
+integer :: i, j
+
+gvar(1)%arr => arr1
+gvar(2)%arr => arr2
+gvar(3)%arr => arr3
+
+gvar(1)%arr = 0
+gvar(2)%arr = 0
+gvar(3)%arr = 0
+
+i = 1
+j = 1
+
+! Here 'gvar(i)' and 'gvar(j)' are the same element, so this should work.
+! This generates a whole-array mapping for gvar(i)%arr, but with the
+! "runtime implicit" bit set so the smaller subarray gvar(j)%arr(1:5) takes
+! precedence.
+
+!$omp target map(gvar(i)%arr, gvar(j)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(j)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(j)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(j)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+! For these ones, we know the array index is the same, so we can just
+! drop the whole-array mapping.
+
+!$omp target map(gvar(i)%arr, gvar(i)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(i)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(i)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(i)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+if (gvar(1)%arr(1).ne.20) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90
new file mode 100644
index 00000000000..5d15808f0da
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90
@@ -0,0 +1,35 @@ 
+! { dg-do run }
+
+type t
+  integer, pointer :: p(:)
+end type t
+
+type(t) :: var(2)
+
+allocate (var(1)%p, source=[1,2,3,5])
+allocate (var(2)%p, source=[2,3,5])
+
+!$omp target map(var(1)%p, var(2)%p)
+var(1)%p(1) = 5
+var(2)%p(2) = 7
+!$omp end target
+
+!$omp target map(var(1)%p(1:3), var(1)%p, var(2)%p)
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(2)%p, var(2)%p(1:3))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(1)%p(1:3), var(2)%p, var(2)%p(2))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+if (var(1)%p(1).ne.8) stop 1
+if (var(2)%p(2).ne.10) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90
new file mode 100644
index 00000000000..9f0edf70890
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90
@@ -0,0 +1,26 @@ 
+! { dg-do run }
+
+type t
+  integer, pointer :: p(:)
+  integer, pointer :: p2(:)
+end type t
+
+type(t) :: var
+integer, target :: tgt(5), tgt2(1000)
+var%p => tgt
+var%p2 => tgt2
+
+p = 0
+p2 = 0
+
+!$omp target map(tgt, tgt2(4:6), var)
+  var%p(1) = 5
+  var%p2(5) = 7
+!$omp end target
+
+if (var%p(1).ne.5) stop 1
+if (var%p2(5).ne.7) stop 2
+
+end
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90
new file mode 100644
index 00000000000..42da7296106
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90
@@ -0,0 +1,29 @@ 
+type t
+integer, pointer :: p2(:)
+end type t
+
+integer, target :: A(5)
+integer, pointer :: p(:), p2(:)
+type(t) :: var
+
+allocate(p2(1:20))
+p => A
+var%p2 => p2
+
+A = 0
+p2 = 0
+
+! These arrays "share original storage", so are unsupported.  This will
+! (correctly) fail with a non-shared address space.
+
+!$omp target map(A(3:4), p2(4:8), p, var%p2)
+A(3) = A(3) + 1
+p2(4) = p2(4) + 2
+!$omp end target
+
+if (A(3).ne.1) stop 1
+if (p2(4).ne.2) stop 2
+
+end program
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90
new file mode 100644
index 00000000000..a47360e10ec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90
@@ -0,0 +1,47 @@ 
+! { dg-do run }
+
+type F
+integer, pointer :: mem(:)
+end type F
+
+type(F) :: fv
+integer, allocatable, target :: arr(:)
+
+allocate(arr(1:20))
+
+fv%mem => arr
+fv%mem = 0
+
+!$omp target enter data map(to: fv%mem(1:10))
+!$omp target map(alloc: fv%mem)
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem(1:10))
+
+if (fv%mem(1).ne.1) stop 1
+
+!$omp target enter data map(to: fv, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv, fv%mem(1:10))
+
+if (fv%mem(1).ne.2) stop 2
+
+!$omp target enter data map(to: fv%mem, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem, fv%mem(1:10))
+
+if (fv%mem(1).ne.3) stop 3
+
+!$omp target enter data map(to: fv%mem)
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem)
+
+if (fv%mem(1).ne.4) stop 4
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray.f90
new file mode 100644
index 00000000000..85f5af3a2a6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray.f90
@@ -0,0 +1,33 @@ 
+! { dg-do run }
+
+program myprog
+type u
+  integer, dimension (:), pointer :: tarr
+end type u
+
+type(u) :: myu
+integer, dimension (12), target :: myarray
+
+myu%tarr => myarray
+
+myu%tarr = 0
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(:))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1:2))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu%tarr)
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+if (myu%tarr(1).ne.4) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90
new file mode 100644
index 00000000000..c7f90131cba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90
@@ -0,0 +1,32 @@ 
+! { dg-do run }
+
+module mymod
+type F
+integer :: a, b, c
+integer, dimension(10) :: d
+end type F
+
+type G
+integer :: x, y
+type(F), pointer :: myf
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+type(F), target :: ftmp
+type(G) :: gvar
+
+gvar%myf => ftmp
+
+gvar%myf%d = 0
+
+!$omp target map(to:gvar%myf) map(tofrom: gvar%myf%b, gvar%myf%d)
+gvar%myf%d(1) = gvar%myf%d(1) + 1
+!$omp end target
+
+if (gvar%myf%d(1).ne.1) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
index 7f3d8174f97..b1d696656c0 100644
--- a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
@@ -36,6 +36,10 @@  program main
   call six ()
   call seven ()
   call eight ()
+  call nine ()
+  call ten ()
+  call eleven ()
+  call twelve ()
 
 contains
   ! Implicitly mapped – but no pointers are mapped
@@ -408,7 +412,180 @@  contains
     !$omp end target
   end subroutine eight
 
-end program main
+  ! This is "subroutine four" but with explicit base-pointer mappings
+  ! (var%f, etc.).
+  subroutine nine()
+    type(t2) :: var
 
-! Fixed by the "Fortran pointers and member mappings" patch
-! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
+    print '(g0)', '==== TESTCASE "nine" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+             uni1 = 4_"abcde", &
+             uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+    allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+!   !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
+!   !$omp&       map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
+    !$omp target map(to: var%f) map(tofrom: var%d(4:7), var%f(2:3), &
+    !$omp&       var%str2(2:3), var%uni2(2:3))
+      if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+
+      if (.not. associated (var%f)) stop 9
+      if (size (var%f) /= 4) stop 10
+      if (any (var%f(2:3) /= [33, 44])) stop 11
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+
+      if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 19
+!     if (.not. associated (var%uni4)) stop 20
+!     if (len (var%uni4) /= 5) stop 21
+!     if (size (var%uni4) /= 2) stop 22
+!     if (var%uni4(2) /= "Go!!!") stop 23
+    !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine nine
+
+  ! This is "subroutine five" but with explicit base-pointer mappings.
+  subroutine ten()
+    type(t2) :: var
+
+    print '(g0)', '==== TESTCASE "ten" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+             uni1 = 4_"abcde", &
+             uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom: var%d(4:7))
+      if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+    !$omp end target
+    !$omp target map(tofrom: var%str2(2:3))
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+    !$omp end target
+
+    !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (any (var%f(2:3) /= [33, 44])) stop 11
+    !$omp end target
+!  !$omp target map(tofrom: var%str4(2:2))
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+!   !$omp end target
+!  !$omp target map(tofrom: var%uni4(2:2))
+!     if (.not. associated (var%uni4)) stop 15
+!     if (len (var%uni4) /= 5) stop 16
+!     if (size (var%uni4) /= 2) stop 17
+!     if (var%uni4(2) /= 4_"Go!!!") stop 18
+!  !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine ten
+
+  ! This is "subroutine six" but with explicit base pointer mappings.
+  subroutine eleven()
+    type(t2) :: var
+
+    print '(g0)', '==== TESTCASE "eleven" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+             uni1 = 4_"abcde", &
+             uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+    allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+!   !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
+!   !$omp                    var%str4(2), var%uni2(3), var%uni4(2))
+    !$omp target map(to: var%f) map(tofrom: var%d(5), var%f(3), &
+    !$omp&                                  var%str2(3), var%uni2(3))
+      if (var%d(5) /= -3*5) stop 4
+      if (var%str2(3) /= "ABCDE") stop 6
+      if (var%uni2(3) /= 4_"ABCDE") stop 7
+
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (var%f(3) /= 44) stop 11
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+!     if (.not. associated (var%uni4)) stop 19
+!     if (len (var%uni4) /= 5) stop 20
+!     if (size (var%uni4) /= 2) stop 21
+!     if (var%uni4(2) /= 4_"Go!!!") stop 22
+    !$omp end target
+
+    deallocate(var%f, var%str4, var%uni4)
+  end subroutine eleven
+
+  ! This is "subroutine seven" but with explicit base-pointer mappings.
+  subroutine twelve()
+    type(t2) :: var
+
+    print '(g0)', '==== TESTCASE "twelve" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+             uni1 = 4_"abcde", &
+             uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+    allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+    !$omp target map(tofrom: var%d(5))
+      if (var%d(5) /= (-3*5)) stop 4
+    !$omp end target
+    !$omp target map(tofrom: var%str2(2:3))
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+    !$omp end target
+    !$omp target map(tofrom: var%uni2(2:3))
+      if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
+    !$omp end target
+
+    !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (any (var%f(2:3) /= [33, 44])) stop 11
+    !$omp end target
+!   !$omp target map(tofrom: var%str4(2:2))
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+!   !$omp end target
+!   !$omp target map(tofrom: var%uni4(2:2))
+!     if (.not. associated (var%uni4)) stop 15
+!     if (len (var%uni4) /= 5) stop 16
+!     if (size (var%uni4) /= 2) stop 17
+!     if (var%uni4(2) /= 4_"Go!!!") stop 18
+!   !$omp end target
+
+    deallocate(var%f, var%str4, var%uni4)
+  end subroutine twelve
+
+end program main