[1/2] OpenMP/Fortran: Map intermediate array descriptor [PR120505]

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 254fc934af1..6bfee916365 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3531,6 +3531,162 @@  get_symbol_rooted_namelist (hash_map<gfc_symbol *,
   return NULL;
 }
 
+/* Helper function for gfc_trans_omp_clauses.  */
+
+static bool
+gfc_map_array_descriptor (
+  tree &node, tree &node2, tree &node3, tree &node4, tree descr, bool openacc,
+  location_t map_loc, stmtblock_t *block, gfc_exec_op op, gfc_omp_namelist *n,
+  hash_map<gfc_symbol *, gfc_omp_namelist *> *&sym_rooted_nl, gfc_se se,
+  gfc_omp_clauses *clauses, bool mid_desc_p)
+{
+  tree type = TREE_TYPE (descr);
+  tree ptr = gfc_conv_descriptor_data_get (descr);
+  ptr = build_fold_indirect_ref (ptr);
+  OMP_CLAUSE_DECL (node) = ptr;
+  int rank = GFC_TYPE_ARRAY_RANK (type);
+  OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, descr, rank);
+  tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+
+  gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (node);
+  if (GOMP_MAP_COPY_TO_P (map_kind) || map_kind == GOMP_MAP_ALLOC)
+    {
+      if (mid_desc_p)
+	{
+	  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+	  OMP_CLAUSE_SIZE (node) = size_int (0);
+	}
+      else
+	map_kind
+	  = ((GOMP_MAP_ALWAYS_P (map_kind) || gfc_expr_attr (n->expr).pointer)
+	       ? GOMP_MAP_ALWAYS_TO
+	       : GOMP_MAP_TO);
+    }
+  else if (n->u.map.op == OMP_MAP_RELEASE || n->u.map.op == OMP_MAP_DELETE)
+    ;
+  else if (op == EXEC_OMP_TARGET_EXIT_DATA || op == EXEC_OACC_EXIT_DATA)
+    map_kind = GOMP_MAP_RELEASE;
+  else
+    map_kind = GOMP_MAP_ALLOC;
+
+  if (!openacc && n->expr->ts.type == BT_CHARACTER && n->expr->ts.deferred)
+    {
+      gcc_assert (se.string_length);
+      tree len = fold_convert (size_type_node, se.string_length);
+      elemsz = gfc_get_char_type (n->expr->ts.kind);
+      elemsz = TYPE_SIZE_UNIT (elemsz);
+      elemsz = fold_build2 (MULT_EXPR, size_type_node, len, elemsz);
+      node4 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
+      OMP_CLAUSE_DECL (node4) = se.string_length;
+      OMP_CLAUSE_SIZE (node4) = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+    }
+  elemsz = fold_convert (gfc_array_index_type, elemsz);
+  OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type,
+					OMP_CLAUSE_SIZE (node), elemsz);
+
+  node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+  if (map_kind == GOMP_MAP_RELEASE || map_kind == GOMP_MAP_DELETE)
+    {
+      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) = descr;
+  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+
+  if (!openacc)
+    {
+      if (n->expr->ts.type == BT_DERIVED
+	  && n->expr->ts.u.derived->attr.alloc_comp)
+	{
+	  /* Save array descriptor for use
+	     in gfc_omp_deep_mapping{,_p,_cnt}; force
+	     evaluate to ensure that it is
+	     not gimplified + is a decl.  */
+	  tree tmp = OMP_CLAUSE_SIZE (node);
+	  tree var = gfc_create_var (TREE_TYPE (tmp), NULL);
+	  gfc_add_modify_loc (map_loc, block, var, tmp);
+	  OMP_CLAUSE_SIZE (node) = var;
+	  gfc_allocate_lang_decl (var);
+	  GFC_DECL_SAVED_DESCRIPTOR (var) = descr;
+	}
+
+      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)
+	return true;
+    }
+
+  if (mid_desc_p && GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (node)))
+    node = NULL_TREE;
+
+  node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH);
+  OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (descr);
+  /* Similar to gfc_trans_omp_array_section (details
+     there), we add/keep the cast for OpenMP to prevent
+     that an 'alloc:' gets added for node3 ('desc.data')
+     as that is part of the whole descriptor (node3).
+     TODO: Remove once the ME handles this properly.  */
+  if (!openacc)
+    OMP_CLAUSE_DECL (node3) = fold_convert (TREE_TYPE (TREE_OPERAND (ptr, 0)),
+					    OMP_CLAUSE_DECL (node3));
+  else
+    STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+  OMP_CLAUSE_SIZE (node3)
+    = size_int (mid_desc_p ? OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT : 0);
+
+  return false;
+}
+
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		       locus where, bool declare_simd = false,
@@ -3544,6 +3700,7 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
   enum omp_clause_code clause_code;
   gfc_omp_namelist *prev = NULL;
   gfc_se se;
+  vec<gfc_symbol *> descriptors = vNULL;
 
   if (clauses == NULL)
     return NULL_TREE;
@@ -4645,6 +4802,8 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		{
 		  gfc_init_se (&se, NULL);
 		  se.expr = gfc_maybe_dereference_var (n->sym, decl);
+		  tree mid_descr = NULL_TREE;
+		  gfc_ref *midref = NULL;
 
 		  for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
 		    {
@@ -4654,6 +4813,12 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    conv_parent_component_references (&se, ref);
 
 			  gfc_conv_component_ref (&se, ref);
+			  if (!mid_descr
+			      && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (se.expr)))
+			    {
+			      mid_descr = se.expr;
+			      midref = ref;
+			    }
 			}
 		      else if (ref->type == REF_ARRAY)
 			{
@@ -4807,156 +4972,11 @@  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 type = TREE_TYPE (inner);
-			  tree ptr = gfc_conv_descriptor_data_get (inner);
-			  ptr = build_fold_indirect_ref (ptr);
-			  OMP_CLAUSE_DECL (node) = ptr;
-			  int rank = GFC_TYPE_ARRAY_RANK (type);
-			  OMP_CLAUSE_SIZE (node)
-			    = gfc_full_array_size (block, inner, rank);
-			  tree elemsz
-			    = TYPE_SIZE_UNIT (gfc_get_element_type (type));
-			  map_kind = OMP_CLAUSE_MAP_KIND (node);
-			  if (GOMP_MAP_COPY_TO_P (map_kind)
-			      || map_kind == GOMP_MAP_ALLOC)
-			    map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
-					 || gfc_expr_attr (n->expr).pointer)
-					? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO);
-			  else if (n->u.map.op == OMP_MAP_RELEASE
-				   || n->u.map.op == OMP_MAP_DELETE)
-			    ;
-			  else if (op == EXEC_OMP_TARGET_EXIT_DATA
-				   || op == EXEC_OACC_EXIT_DATA)
-			    map_kind = GOMP_MAP_RELEASE;
-			  else
-			    map_kind = GOMP_MAP_ALLOC;
-			  if (!openacc
-			      && n->expr->ts.type == BT_CHARACTER
-			      && n->expr->ts.deferred)
-			    {
-			      gcc_assert (se.string_length);
-			      tree len = fold_convert (size_type_node,
-						       se.string_length);
-			      elemsz = gfc_get_char_type (n->expr->ts.kind);
-			      elemsz = TYPE_SIZE_UNIT (elemsz);
-			      elemsz = fold_build2 (MULT_EXPR, size_type_node,
-						    len, elemsz);
-			      node4 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
-			      OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
-			      OMP_CLAUSE_DECL (node4) = se.string_length;
-			      OMP_CLAUSE_SIZE (node4)
-				= TYPE_SIZE_UNIT (gfc_charlen_type_node);
-			    }
-			  elemsz = fold_convert (gfc_array_index_type, elemsz);
-			  OMP_CLAUSE_SIZE (node)
-			    = fold_build2 (MULT_EXPR, gfc_array_index_type,
-					   OMP_CLAUSE_SIZE (node), elemsz);
-			  node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
-			  if (map_kind == GOMP_MAP_RELEASE
-			      || map_kind == GOMP_MAP_DELETE)
-			    {
-			      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)
-			    {
-			      if (n->expr->ts.type == BT_DERIVED
-				  && n->expr->ts.u.derived->attr.alloc_comp)
-				{
-				  /* Save array descriptor for use
-				     in gfc_omp_deep_mapping{,_p,_cnt}; force
-				     evaluate to ensure that it is
-				     not gimplified + is a decl.  */
-				  tree tmp = OMP_CLAUSE_SIZE (node);
-				  tree var = gfc_create_var (TREE_TYPE (tmp),
-							     NULL);
-				  gfc_add_modify_loc (map_loc, block,
-						      var, tmp);
-				  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;
-			    }
-			  node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node3,
-						   GOMP_MAP_ATTACH_DETACH);
-			  OMP_CLAUSE_DECL (node3)
-			    = gfc_conv_descriptor_data_get (inner);
-			  /* Similar to gfc_trans_omp_array_section (details
-			     there), we add/keep the cast for OpenMP to prevent
-			     that an 'alloc:' gets added for node3 ('desc.data')
-			     as that is part of the whole descriptor (node3).
-			     TODO: Remove once the ME handles this properly.  */
-			  if (!openacc)
-			    OMP_CLAUSE_DECL (node3)
-				= fold_convert (TREE_TYPE (TREE_OPERAND(ptr, 0)),
-						OMP_CLAUSE_DECL (node3));
-			  else
-			    STRIP_NOPS (OMP_CLAUSE_DECL (node3));
-			  OMP_CLAUSE_SIZE (node3) = size_int (0);
+			  bool drop_mapping = gfc_map_array_descriptor (
+			    node, node2, node3, node4, inner, openacc, map_loc,
+			    block, op, n, sym_rooted_nl, se, clauses, false);
+			  if (drop_mapping)
+			    continue;
 			}
 		      else
 			OMP_CLAUSE_DECL (node) = inner;
@@ -4972,6 +4992,31 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		    }
 		  else
 		    gcc_unreachable ();
+
+		  /* Map intermediate array descriptor.  */
+		  if (!openacc && mid_descr != NULL_TREE && mid_descr != inner
+		      && !descriptors.contains (midref->u.c.sym))
+		    {
+		      descriptors.safe_push (midref->u.c.sym);
+
+		      tree node1 = copy_node (node);
+		      tree node2 = NULL_TREE;
+		      tree node3 = NULL_TREE;
+		      tree node4 = NULL_TREE;
+		      gfc_map_array_descriptor (node1, node2, node3, node4,
+						mid_descr, openacc, map_loc,
+						block, op, n, sym_rooted_nl, se,
+						clauses, true);
+
+		      if (node1 != NULL_TREE)
+			omp_clauses = gfc_trans_add_clause (node1, omp_clauses);
+		      if (node2 != NULL_TREE)
+			omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
+		      if (node3 != NULL_TREE)
+			omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
+		      if (node4 != NULL_TREE)
+			omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+		    }
 		}
 	      else
 		sorry_at (gfc_get_location (&n->where), "unhandled expression");
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 751b4697271..8559b7ff2ad 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11511,9 +11511,13 @@  omp_mapped_by_containing_struct (hash_map<tree_operand_hash_no_se,
 	}
       if (wholestruct)
 	{
+	  tree desc = OMP_CLAUSE_CHAIN (*(*wholestruct)->grp_start);
+	  if (desc != NULL_TREE && omp_map_clause_descriptor_p (desc))
+	    goto next;
 	  *mapped_by_group = *wholestruct;
 	  return true;
 	}
+    next:
       decl = wsdecl;
     }
 
@@ -13390,6 +13394,35 @@  omp_build_struct_sibling_lists (enum tree_code code,
       tail = added_tail;
     }
 
+  /* Find each attach node whose bias needs to be adjusted and move it to the
+   * group containing its pointee, right after the struct node.  */
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      tree c = *grp->grp_start;
+      if (c != NULL && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+	  && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_TO_PSET
+	  && OMP_CLAUSE_MAP_KIND (grp->grp_end) == GOMP_MAP_ATTACH_DETACH
+	  && OMP_CLAUSE_SIZE (grp->grp_end)
+	       == size_int (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT))
+	{
+	  tree *cp;
+	  for (cp = &OMP_CLAUSE_CHAIN (c); cp != NULL;
+	       cp = &OMP_CLAUSE_CHAIN (*cp))
+	    if (*cp == grp->grp_end)
+	      {
+		c = *cp;
+		break;
+	      }
+
+	  tree base = OMP_CLAUSE_DECL (c);
+	  gcc_assert (TREE_CODE (base) == NOP_EXPR);
+	  base = build_fold_indirect_ref (base);
+	  tree *struct_node = struct_map_to_clause->get (base);
+	  omp_siblist_move_node_after (c, cp, &OMP_CLAUSE_CHAIN (*struct_node));
+	}
+    }
+
   /* Now we have finished building the struct sibling lists, reprocess
      newly-added "attach" nodes: we need the address of the first
      mapped element of each struct sibling list for the bias of the attach
@@ -13416,7 +13449,9 @@  omp_build_struct_sibling_lists (enum tree_code code,
 	   base they attach to).  We should only have created the
 	   ATTACH_DETACH node after GOMP_MAP_STRUCT for a target region, so
 	   this should never be true.  */
-	gcc_assert ((region_type & ORT_TARGET) != 0);
+	// This is no longer true. See zlas in gomp_map_vars_internal
+	// (libgomp/target.c).
+	// gcc_assert ((region_type & ORT_TARGET) != 0);
 
 	/* This is the first sorted node in the struct sibling list.  Use it
 	   to recalculate the correct bias to use.
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
new file mode 100644
index 00000000000..b2bec2501fa
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
@@ -0,0 +1,49 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! PR fortran/120505
+
+! Check that mapping nested allocatable DT components triggers required
+! additional mappings for the outer array descriptor.
+
+module m
+
+  type field_type
+    real(kind=8), allocatable :: density0(:,:), density1(:,:)
+  end type field_type
+
+  type tile_type
+    type(field_type) :: field
+  end type tile_type
+
+  type chunk_type
+    real(kind=8), allocatable :: left_rcv_buffer(:)
+    type(tile_type), allocatable :: tiles(:)
+  end type chunk_type
+
+  type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2])
+chunk%tiles(1)%field%density1 = reshape([5,6,7,8],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp   map(to: chunk%tiles(1)%field%density0) &
+!$omp   map(to: chunk%tiles(1)%field%density1) &
+!$omp   map(to: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { #pragma omp target enter data map\(alloc:\*\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[len: 0\] \[runtime_implicit\]\) map\(to:chunk\.tiles \[pointer set, len: 64\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: [0-9]+\]\) } 1 "original" } }
+
+!$omp target exit data &
+!$omp   map(from: chunk%tiles(1)%field%density0) &
+!$omp   map(from: chunk%tiles(1)%field%density1) &
+!$omp   map(from: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { #pragma omp target exit data map\(release:chunk\.tiles \[pointer set, len: 64\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: [0-9]+\]\) } 1 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
new file mode 100644
index 00000000000..a44c4d72960
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
@@ -0,0 +1,47 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+! PR fortran/120505
+
+! Check that the bias into the inner derived type is correctly computed on 
+! target enter data. For target exit data, the bias is ignored so just check
+! that detach is present.
+
+module m
+
+  type field_type
+    real(kind=8), allocatable :: density0(:,:), density1(:,:)
+  end type field_type
+
+  type tile_type
+    type(field_type) :: field
+  end type tile_type
+
+  type chunk_type
+    real(kind=8), allocatable :: left_rcv_buffer(:)
+    type(tile_type), allocatable :: tiles(:)
+  end type chunk_type
+
+  type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp   map(to: chunk%tiles(1)%field%density1) &
+!$omp   map(to: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { map\(struct_unord:MEM <struct tile_type\[0:\]> \[\(struct tile_type\[0:\] \*\)_[0-9]+\] \[len: 1\]\) map\(to:MEM <struct tile_type\[0:\]> \[\(struct tile_type\[0:\] \*\)_[0-9]+\]\[_[0-9]+\]\.field\.density1 \[pointer set, len: 88\]\) map\(attach:chunk\.tiles\.data \[bias: _[0-9]+\]\) } 1 "gimple" } }
+
+!$omp target exit data &
+!$omp   map(from: chunk%tiles(1)%field%density1) &
+!$omp   map(from: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { map\(release:chunk\.tiles \[pointer set, len: 64\]\) map\(detach:chunk\.tiles\.data \[bias: [0-9]+\]\)} 1 "gimple" } }
+
+end
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 7bdb474a253..555b54b7ced 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -667,6 +667,7 @@  enum omp_clause_fallback_kind {
   OMP_CLAUSE_FALLBACK_NULL
 };
 
+#define OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT -1
 
 /* memory-order-clause on OpenMP atomic/flush constructs or
    argument of atomic_default_mem_order clause.  */
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
new file mode 100644
index 00000000000..d81a31491bb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
@@ -0,0 +1,56 @@ 
+! { dg-do run }
+
+! PR fortran/120505
+
+! Check that mapping nested allocatable DT components triggers required
+! additional mappings for the outer array descriptor.
+
+module m
+
+  type field_type
+    real(kind=8), allocatable :: density0(:,:), density1(:,:)
+  end type field_type
+
+  type tile_type
+    type(field_type) :: field
+  end type tile_type
+
+  type chunk_type
+    real(kind=8), allocatable :: left_rcv_buffer(:)
+    type(tile_type), allocatable :: tiles(:)
+  end type chunk_type
+
+  type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2])
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp   map(to: chunk%tiles(1)%field%density0) &
+!$omp   map(to: chunk%tiles(1)%field%density1) &
+!$omp   map(to: chunk%left_rcv_buffer)
+
+!$omp target
+  if (any (chunk%tiles(1)%field%density0 /= reshape([1,2,3,4],[2,2]))) stop 1
+  if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+  chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 + 7
+  chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5
+  chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp   map(from: chunk%tiles(1)%field%density0) &
+!$omp   map(from: chunk%tiles(1)%field%density1) &
+!$omp   map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field%density0 /= 7 + reshape([1,2,3,4],[2,2]))) stop 1
+if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
new file mode 100644
index 00000000000..13fad59415b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
@@ -0,0 +1,50 @@ 
+! { dg-do run }
+
+! PR fortran/120505
+
+! Check that a nested allocatable DT component is mapped properly even when the
+! first component is *not* mapped.
+
+module m
+
+  type field_type
+    real(kind=8), allocatable :: density0(:,:), density1(:,:)
+  end type field_type
+
+  type tile_type
+    type(field_type) :: field
+  end type tile_type
+
+  type chunk_type
+    real(kind=8), allocatable :: left_rcv_buffer(:)
+    type(tile_type), allocatable :: tiles(:)
+  end type chunk_type
+
+  type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp   map(to: chunk%tiles(1)%field%density1) &
+!$omp   map(to: chunk%left_rcv_buffer)
+
+!$omp target
+  if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+  chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5
+  chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp   map(from: chunk%tiles(1)%field%density1) &
+!$omp   map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1
+
+end