Re: [PATCH v7 3/5] OpenMP: Pointers and member mappings

[Julian Brown <julian@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 5270 bytes --]

Hi Tobias!

On Wed, 6 Dec 2023 12:36:34 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:

> Hi Julian,
> 
> LGTM, except for:
> 
> * The 'target exit data' handling - comments below - looks a bit
> fishy/inconsistent.
> 
> I intent to have a closer look with more code context, but maybe you
> should have a look at it as well.
> 
> BTW: Fortran deep-mapping is not yet on mainline. Are you aware of
> changes or in particular testcases on OG13 related to your patch
> series that should be included when upstreaming that auto-mapping of
> allocatable components patch?

I thought I'd adjusted some tests to use "pointer" instead of
"allocatable" at some point for mainline submission, but now I can't
find where I'm thinking of.  So possibly.  (I'll keep an eye out...)

> > +      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
> > +       || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
> > +       || op == EXEC_OMP_TARGET_EXIT_DATA)
> >       {
> > [...]
> > +       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;  
> 
> For '!$omp target exit data map(delete: array)' this looks wrong as it
> replaces 'delete' by 'release' for the descriptor - while for
>   '!$omp target (data) map(delete: array)'
> it remains 'delete'.
> 
> Thus, I wonder whether that shouldn't be instead
>    OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
>    ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE;

I've fixed that as you suggest.  Actually I've made OpenACC use the new
node layout as well, since (a) it works and (b) it was weirdly
inconsistent before.  That is, exit data directives will no longer use
e.g.:

  GOMP_MAP_FROM
  GOMP_MAP_TO_PSET
  GOMP_MAP_ATTACH_DETACH

but instead,

  GOMP_MAP_FROM
  GOMP_MAP_RELEASE (with OMP_CLAUSE_RELEASE_DESCRIPTOR set)
  GOMP_MAP_ATTACH_DETACH

actually the current state is that GOMP_MAP_TO_PSET will be used for
the descriptor on an "exit data" directive if you refer to the whole
array, but GOMP_MAP_RELEASE (etc.) will be used if you refer to an array
section (without the flag newly added in this patch, of course). I
don't think there's any reason to maintain that inconsistency.

> We later have the following; just reading the patch, I wonder whether
> GOMP_TO_PSET is correct for a generic 'target exit data' here or not.
>  It seems at a glance as if an "|| op == 'target exit data'" is
> missing here:
> 
> > -                       if (openacc)
> > -                         OMP_CLAUSE_SET_MAP_KIND (desc_node,
> > +                       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);  
> 
> (Here, without setting OMP_CLAUSE_RELEASE_DESCRIPTOR.)
> 
> And for completeness, we later have:
> 
> > +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;  

In that section, the "exit data" directive case is handled a few lines
higher up the function, with a condition that (now) reads:

  else if (op == EXEC_OMP_TARGET_EXIT_DATA
           || op == EXEC_OACC_EXIT_DATA)
    map_kind = GOMP_MAP_RELEASE;

so I think we're OK -- unless I missed something?  (Removing the
OpenACC special case and inverting a conditional simplifies the logic
here a bit.)

> > +                               /* Save array descriptor for use
> > +                                  in
> > gfc_omp_deep_mapping{,_p,_cnt}; force
> > +                                  evaluate to ensure that it is
> > +                                  not gimplified + is a decl.  */  
> This is part of my Fortran allocatable-components deep-mapping patch
> that is currently only on OG9 (?) to OG13, badly needs to be
> upstreamed but required that Jakub had a look at it - well, I still
> would like that he has a look at the omp-low.cc parts and it needs to
> be re-diffed.
> 
> Hence, while I wouldn't mind to keep it to avoid more diffing work, I
> think it will be cleaner not to keep it here.

Agreed -- that looks like a bad bit of merge resolution on my part, in
fact.  I've removed it from this patch.

I've re-tested this version.  Does it look better now?

Thanks,

Julian


[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: pointers-and-member-mappings-8.diff --]
[-- Type: text/x-patch, Size: 73944 bytes --]

commit 86563a13497a324d4e72790551706bb5c2c22f76
Author: Julian Brown <julian@codesourcery.com>
Date:   Mon Oct 17 16:44:31 2022 +0000

    OpenMP: Pointers and member mappings
    
    This patch changes the mapping node arrangement used for array components
    of derived types in order to accommodate for changes made in the previous
    patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed
    derived-type members instead of "GOMP_MAP_ALWAYS_POINTER".
    
    We change the mapping nodes used for a derived-type mapping like this:
    
      type T
      integer, pointer, dimension(:) :: arrptr
      end type T
    
      type(T) :: tvar
      [...]
      !$omp target map(tofrom: tvar%arrptr)
    
    So that the nodes used look like this:
    
      1) map(to: tvar%arrptr)   -->
      GOMP_MAP_TO [implicit]  *tvar%arrptr%data  (the array data)
      GOMP_MAP_TO_PSET        tvar%arrptr        (the descriptor)
      GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data
    
      2) map(tofrom: tvar%arrptr(3:8)   -->
      GOMP_MAP_TOFROM         *tvar%arrptr%data(3)  (size 8-3+1, etc.)
      GOMP_MAP_TO_PSET        tvar%arrptr
      GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data      (bias 3, etc.)
    
    In this case, we can determine in the front-end that the
    whole-array/pointer mapping (1) is only needed to map the pointer
    -- so we drop it entirely.  (Note also that we set -- early -- the
    OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P flag for whole-array-via-pointer
    mappings. See below.)
    
    In the middle end, we process mappings using the struct sibling-list
    handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle
    of the group of three mapping nodes to the proper sorted position after
    the GOMP_MAP_STRUCT mapping:
    
      GOMP_MAP_STRUCT   tvar     (len: 1)
      GOMP_MAP_TO_PSET  tvar%arr (size: 64, etc.)  <--. moved here
      [...]                                           |
      GOMP_MAP_TOFROM         *tvar%arrptr%data(3) ___|
      GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data
    
    In another case, if we have an array of derived-type values "dtarr",
    and mappings like:
    
      i = 1
      j = 1
      map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8))
    
    We still map the same way, but this time we cannot prove that the base
    expressions "dtarr(i) and "dtarr(j)" are the same in the front-end.
    So we keep both mappings, but we move the "[implicit]" mapping of the
    full-array reference to the end of the clause list in gimplify.cc (by
    adjusting the topological sorting algorithm):
    
      GOMP_MAP_STRUCT         dtvar  (len: 2)
      GOMP_MAP_TO_PSET        dtvar(i)%arrptr
      GOMP_MAP_TO_PSET        dtvar(j)%arrptr
      [...]
      GOMP_MAP_TOFROM         *dtvar(j)%arrptr%data(3)  (size: 8-3+1)
      GOMP_MAP_ATTACH_DETACH  dtvar(j)%arrptr%data
      GOMP_MAP_TO [implicit]  *dtvar(i)%arrptr%data(1)  (size: whole array)
      GOMP_MAP_ATTACH_DETACH  dtvar(i)%arrptr%data
    
    Always moving "[implicit]" full-array mappings after array-section
    mappings (without that bit set) means that we'll avoid copying the whole
    array unnecessarily -- even in cases where we can't prove that the arrays
    are the same.
    
    The patch also fixes some bugs with "enter data" and "exit data"
    directives with this new mapping arrangement.  Also now if you have
    mappings like this:
    
      #pragma omp target enter data map(to: dv, dv%arr(1:20))
    
    The whole of the derived-type variable "dv" is mapped, so the
    GOMP_MAP_TO_PSET for the array-section mapping can be dropped:
    
      GOMP_MAP_TO            dv
    
      GOMP_MAP_TO            *dv%arr%data
      GOMP_MAP_TO_PSET       dv%arr <-- deleted (array section mapping)
      GOMP_MAP_ATTACH_DETACH dv%arr%data
    
    To accommodate for recent changes to mapping nodes made by
    Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET
    for "exit data" directives, in favour of using the "correct"
    GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion.  A new
    flag is introduced so the middle-end knows when the latter two kinds
    are being used specifically for an array descriptor.
    
    This version of the patch fixes "omp target exit data" handling
    for GOMP_MAP_DELETE, and adds pretty-printing dump output
    for the OMP_CLAUSE_RELEASE_DESCRIPTOR flag (for a little extra
    clarity).
    
    Also I noticed the handling of descriptors on *OpenACC*
    exit-data directives was inconsistent, so I've made those use
    GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the new flag in the same way as
    OpenMP too.  In the end it doesn't actually matter to the runtime,
    which handles GOMP_MAP_RELEASE/GOMP_MAP_DELETE/GOMP_MAP_TO_PSET for
    array descriptors on OpenACC "exit data" directives the same, anyway,
    and doing it this way in the FE avoids needless divergence.
    
    I've added a couple of new tests (gomp/target-enter-exit-data.f90 and
    goacc/enter-exit-data-2.f90).
    
    2023-12-07  Julian Brown  <julian@codesourcery.com>
    
    gcc/fortran/
            * dependency.cc (gfc_omp_expr_prefix_same): New function.
            * dependency.h (gfc_omp_expr_prefix_same): Add prototype.
            * gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2"
            union.
            * trans-openmp.cc (dependency.h): Include.
            (gfc_trans_omp_array_section): Adjust mapping node arrangement for
            array descriptors.  Use GOMP_MAP_TO_PSET or
            GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DESCRIPTOR
            flag set.
            (gfc_symbol_rooted_namelist): New function.
            (gfc_trans_omp_clauses): Check subcomponent and subarray/element
            accesses elsewhere in the clause list for pointers to derived types or
            array descriptors, and adjust or drop mapping nodes appropriately.
            Adjust for changes to mapping node arrangement.
            (gfc_trans_oacc_executable_directive): Pass code op through.
    
    gcc/
            * gimplify.cc (omp_map_clause_descriptor_p): New function.
            (build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base): Use
            above function.
            (omp_tsort_mapping_groups): Process nodes that have
            OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't.  Add
            enter_exit_data parameter.
            (omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if
            we're mapping the whole containing derived-type variable.
            (omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling.
            Remove GOMP_MAP_ALWAYS_POINTER handling.
            (gimplify_scan_omp_clauses): Pass enter_exit argument to
            omp_tsort_mapping_groups.  Don't adjust/remove GOMP_MAP_TO_PSET
            mappings for derived-type components here.
            * tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro.
            * tree-pretty-print.cc (dump_omp_clause): Show
            OMP_CLAUSE_RELEASE_DESCRIPTOR in dump output (with
            GOMP_MAP_TO_PSET-like syntax).
    
    gcc/testsuite/
            * gfortran.dg/goacc/enter-exit-data-2.f90: New test.
            * gfortran.dg/goacc/finalize-1.f: Adjust scan output.
            * gfortran.dg/gomp/map-9.f90: Adjust scan output.
            * gfortran.dg/gomp/map-subarray-2.f90: New test.
            * gfortran.dg/gomp/map-subarray.f90: New test.
            * gfortran.dg/gomp/target-enter-exit-data.f90: New test.
    
    libgomp/
            * testsuite/libgomp.fortran/map-subarray.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-2.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-3.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-4.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-6.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-7.f90: New test.
            * testsuite/libgomp.fortran/map-subarray-8.f90: New test.
            * testsuite/libgomp.fortran/map-subcomponents.f90: New test.
            * testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for
            descriptor-mapping changes.  Remove XFAIL.

diff --git a/gcc/fortran/dependency.cc b/gcc/fortran/dependency.cc
index 632b3985723..94647032ab5 100644
--- a/gcc/fortran/dependency.cc
+++ b/gcc/fortran/dependency.cc
@@ -2337,3 +2337,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 fbbede8b22c..a2458f79034 100644
--- a/gcc/fortran/dependency.h
+++ b/gcc/fortran/dependency.h
@@ -40,5 +40,6 @@ int gfc_expr_is_one (gfc_expr *, int);
 bool gfc_dep_resolver (gfc_ref *, gfc_ref *, gfc_reverse *,
 		      bool identical = false);
 bool 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 aa3f6cb70b4..6be0f853253 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1380,6 +1380,7 @@ typedef struct gfc_omp_namelist
       gfc_namespace *ns;
       gfc_expr *allocator;
       struct gfc_symbol *traits_sym;
+      struct gfc_omp_namelist *duplicate_of;
     } u2;
   struct gfc_omp_namelist *next;
   locus where;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 33216f5c7f2..b9ab8860bd1 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -40,6 +40,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "omp-general.h"
 #include "omp-low.h"
 #include "memmodel.h"  /* For MEMMODEL_ enums.  */
+#include "dependency.h"
 
 #undef GCC_DIAG_STYLE
 #define GCC_DIAG_STYLE __gcc_tdiag__
@@ -2491,36 +2492,24 @@ 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
+	  || op == EXEC_OACC_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
+	    = OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE ? GOMP_MAP_DELETE
+							    : GOMP_MAP_RELEASE;
+	  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);
@@ -2624,6 +2613,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,
@@ -2641,6 +2697,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];
@@ -3664,6 +3722,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)
@@ -3719,7 +3825,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);
@@ -3738,7 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  else if (n->u.map_op == OMP_MAP_RELEASE
 				   || n->u.map_op == OMP_MAP_DELETE)
 			    ;
-			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+			  else if (op == EXEC_OMP_TARGET_EXIT_DATA
+				   || op == EXEC_OACC_EXIT_DATA)
 			    map_kind = GOMP_MAP_RELEASE;
 			  else
 			    map_kind = GOMP_MAP_ALLOC;
@@ -3764,24 +3870,78 @@ 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,
-						     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
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  if (map_kind == GOMP_MAP_RELEASE
+			      || map_kind == GOMP_MAP_DELETE)
 			    {
-			      node2 = node;
-			      node = desc_node;  /* Put first.  */
+			      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+			      OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;
+			    }
+			  else
+			    OMP_CLAUSE_SET_MAP_KIND (node2,
+						     GOMP_MAP_TO_PSET);
+			  OMP_CLAUSE_DECL (node2) = inner;
+			  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+			  if (!openacc)
+			    {
+			      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,
@@ -3945,6 +4105,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;
@@ -4787,7 +4964,7 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
 
   gfc_start_block (&block);
   oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-					code->loc, false, true);
+					code->loc, 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 1e34bde19be..57e3252608a 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9146,6 +9146,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
@@ -9181,9 +9200,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);
@@ -9369,7 +9386,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);
@@ -9464,7 +9481,7 @@ omp_group_last (tree *start_p)
 		     == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
-		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+		 || omp_map_clause_descriptor_p (nc)))
 	{
 	  tree nc2 = OMP_CLAUSE_CHAIN (nc);
 	  if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
@@ -9631,33 +9648,32 @@ 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:
-	  case GOMP_MAP_DETACH:
-	    return *grp->grp_start;
+	case GOMP_MAP_ALWAYS_POINTER:
+	case GOMP_MAP_ATTACH_DETACH:
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	case GOMP_MAP_DETACH:
+	  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:
@@ -10005,18 +10021,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;
@@ -10419,6 +10462,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.
 
@@ -10598,6 +10646,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)))
@@ -10996,11 +11055,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)
     {
@@ -11023,7 +11088,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
@@ -11254,7 +11330,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
@@ -11310,7 +11386,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;
@@ -11326,7 +11402,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
@@ -11359,7 +11437,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,
@@ -11504,8 +11592,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;
 	}
@@ -11781,6 +11868,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
@@ -11795,7 +11884,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/goacc/enter-exit-data-2.f90 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data-2.f90
new file mode 100644
index 00000000000..6a16c8a89ea
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data-2.f90
@@ -0,0 +1,38 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+type t
+integer, pointer :: arr(:)
+end type t
+
+type(t) :: var
+
+allocate (var%arr(1:100))
+
+!$acc enter data copyin(var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+!$acc exit data delete(var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+
+!$acc enter data create(var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+!$acc exit data finalize delete(var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\) finalize;$} 1 "original" } }
+
+
+!$acc enter data copyin(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+!$acc exit data delete(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+
+!$acc enter data create(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+!$acc exit data finalize delete(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\) finalize;$} 1 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 23f0ffc627e..63beb479453 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -20,8 +20,8 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -32,6 +32,6 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
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..26e113f4fdd
--- /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\(release:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[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 \[pointer set, 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 \[pointer set, 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/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90 b/gcc/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90
new file mode 100644
index 00000000000..c14a11dacd9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90
@@ -0,0 +1,39 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+type t
+integer, pointer :: arr(:)
+end type t
+
+type(t) :: var
+
+allocate (var%arr(1:100))
+
+!$omp target enter data map(to: var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+!$omp target exit data map(release: var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(alloc: var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+!$omp target exit data map(delete: var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(delete:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(delete:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(to: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+!$omp target exit data map(release: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(alloc: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+!$omp target exit data map(delete: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(delete:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(delete:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+
+end
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 1fadd752d05..6e1c386cb02 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1050,6 +1050,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	    case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	      pp_string (pp, " [bias: ");
 	      break;
+	    case GOMP_MAP_RELEASE:
+	    case GOMP_MAP_DELETE:
+	      if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+		  && OMP_CLAUSE_RELEASE_DESCRIPTOR (clause))
+		{
+		  pp_string (pp, " [pointer set, len: ");
+		  break;
+		}
+	      /* Fallthrough.  */
 	    default:
 	      pp_string (pp, " [len: ");
 	      break;
diff --git a/gcc/tree.h b/gcc/tree.h
index 4fd75e6131b..2b69d84d0a6 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1829,6 +1829,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 a release/delete node which refers to a (Fortran) array
+   descriptor.  */
+#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