Fortran/OpenMP: Fix mapping of array descriptors and deferred-length strings

Previously, array descriptors might have been mapped as 'alloc'
instead of 'to' for 'alloc', not updating the array bounds. The
'alloc' could also appear for 'data exit', failing with a libgomp
assert. In some cases, either array descriptors or deferred-length
string's length variable was not mapped. And, finally, some offset
calculations with array-sections mappings went wrong.

Additionally, the patch now unmaps for scalar allocatables/pointers
the GOMP_MAP_POINTER, avoiding stale mappings.

The testcases contain some comment-out tests which require follow-up
work and for which PR exist. Those mostly relate to deferred-length
strings which have several issues beyong OpenMP support.

gcc/fortran/ChangeLog:

	* trans-decl.cc (gfc_get_symbol_decl): Add attributes
	such as 'declare target' also to hidden artificial
	variable for deferred-length character variables.
	* trans-openmp.cc (gfc_trans_omp_array_section,
	gfc_trans_omp_clauses, gfc_trans_omp_target_exit_data):
	Improve mapping of array descriptors and deferred-length
	string variables.

gcc/ChangeLog:

	* gimplify.cc (gimplify_scan_omp_clauses): Remove Fortran
	special case.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/target-enter-data-3.f90: Uncomment
	'target exit data'.
	* testsuite/libgomp.fortran/target-enter-data-4.f90: New test.
	* testsuite/libgomp.fortran/target-enter-data-5.f90: New test.
	* testsuite/libgomp.fortran/target-enter-data-6.f90: New test.
	* testsuite/libgomp.fortran/target-enter-data-7.f90: New test.

gcc/testsuite/
	* gfortran.dg/goacc/finalize-1.f: Update dg-tree; shows a fix
	for 'finalize' as a ptr is now 'delete' instead of 'release'.
	* gfortran.dg/gomp/pr78260-2.f90: Likewise as elem-size calc moved
	to if (allocated) block
	* gfortran.dg/gomp/target-exit-data.f90: Likewise as a var is now a
	replaced by a MEM< _25 > expression.
	* gfortran.dg/gomp/map-9.f90: Update dg-scan-tree-dump.
	* gfortran.dg/gomp/map-10.f90: New test.

13 files changed, 1877 insertions, 140 deletions

diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index 18e2b8b78b4..55a13ee1b3d 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -1827,6 +1827,8 @@ gfc_get_symbol_decl (gfc_symbol * sym)
   /* Add attributes to variables.  Functions are handled elsewhere.  */
   attributes = add_attributes_to_decl (sym->attr, NULL_TREE);
   decl_attributes (&decl, attributes, 0);
+  if (sym->ts.deferred && VAR_P (length))
+    decl_attributes (&length, attributes, 0);
 
   /* Symbols from modules should have their assembler names mangled.
      This is done here rather than in gfc_finish_var_decl because it
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 96aecdd1cb3..9b6ff939128 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2403,33 +2403,50 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
 /* Translate an array section or array element.  */
 
 static void
-gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
-			     tree decl, bool element, gomp_map_kind ptr_kind,
-			     tree &node, tree &node2, tree &node3, tree &node4)
+gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
+			     gfc_omp_namelist *n, tree decl, bool element,
+			     gomp_map_kind ptr_kind, tree &node, tree &node2,
+			     tree &node3, tree &node4)
 {
   gfc_se se;
   tree ptr, ptr2;
   tree elemsz = NULL_TREE;
 
   gfc_init_se (&se, NULL);
-
   if (element)
     {
       gfc_conv_expr_reference (&se, n->expr);
       gfc_add_block_to_block (block, &se.pre);
       ptr = se.expr;
-      OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
-      elemsz = OMP_CLAUSE_SIZE (node);
     }
   else
     {
       gfc_conv_expr_descriptor (&se, n->expr);
       ptr = gfc_conv_array_data (se.expr);
+    }
+  if (n->expr->ts.type == BT_CHARACTER && n->expr->ts.deferred)
+    {
+      gcc_assert (se.string_length);
+      tree len = gfc_evaluate_now (se.string_length, block);
+      elemsz = gfc_get_char_type (n->expr->ts.kind);
+      elemsz = TYPE_SIZE_UNIT (elemsz);
+      elemsz = fold_build2 (MULT_EXPR, size_type_node,
+			    fold_convert (size_type_node, len), elemsz);
+    }
+  if (element)
+    {
+      if (!elemsz)
+	elemsz = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
+      OMP_CLAUSE_SIZE (node) = elemsz;
+    }
+  else
+    {
       tree type = TREE_TYPE (se.expr);
       gfc_add_block_to_block (block, &se.pre);
       OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
 						    GFC_TYPE_ARRAY_RANK (type));
-      elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+      if (!elemsz)
+	elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
       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);
@@ -2441,7 +2458,11 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
 
   if (POINTER_TYPE_P (TREE_TYPE (decl))
       && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
-      && ptr_kind == GOMP_MAP_POINTER)
+      && ptr_kind == GOMP_MAP_POINTER
+      && op != EXEC_OMP_TARGET_EXIT_DATA
+      && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_RELEASE
+      && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_DELETE)
+
     {
       node4 = build_omp_clause (input_location,
 				OMP_CLAUSE_MAP);
@@ -2455,13 +2476,13 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
 	   && n->expr->ts.deferred)
     {
       gomp_map_kind map_kind;
-      if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
-	map_kind = GOMP_MAP_TO;
-      else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
-	       || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
 	map_kind = OMP_CLAUSE_MAP_KIND (node);
+      else if (op == EXEC_OMP_TARGET_EXIT_DATA
+	       || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE)
+	map_kind = GOMP_MAP_RELEASE;
       else
-	map_kind = GOMP_MAP_ALLOC;
+	map_kind = GOMP_MAP_TO;
       gcc_assert (se.string_length);
       node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
@@ -2476,7 +2497,18 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
       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 (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
+      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
+	{
+	  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;
@@ -2487,11 +2519,11 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
 	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
 	  node2 = desc_node;
 	}
-      node3 = build_omp_clause (input_location,
-				OMP_CLAUSE_MAP);
+      if (op == EXEC_OMP_TARGET_EXIT_DATA)
+	return;
+      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);
+      OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
       /* This purposely does not include GOMP_MAP_ALWAYS_POINTER.  The extra
 	 cast prevents gimplify.cc from recognising it as being part of the
 	 struct - and adding an 'alloc: for the 'desc.data' pointer, which
@@ -2595,7 +2627,7 @@ handle_iterator (gfc_namespace *ns, stmtblock_t *iter_block, tree block)
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		       locus where, bool declare_simd = false,
-		       bool openacc = false)
+		       bool openacc = false, gfc_exec_op op = EXEC_NOP)
 {
   tree omp_clauses = NULL_TREE, prev_clauses, chunk_size, c;
   tree iterator = NULL_TREE;
@@ -3026,6 +3058,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      tree node2 = NULL_TREE;
 	      tree node3 = NULL_TREE;
 	      tree node4 = NULL_TREE;
+	      tree node5 = NULL_TREE;
 
 	      /* OpenMP: automatically map pointer targets with the pointer;
 		 hence, always update the descriptor/pointer itself.  */
@@ -3130,6 +3163,24 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  || (n->expr->ref->type == REF_ARRAY
 		      && n->expr->ref->u.ar.type == AR_FULL))
 		{
+		  gomp_map_kind map_kind;
+		  tree type = TREE_TYPE (decl);
+		  if (n->sym->ts.type == BT_CHARACTER
+		      && n->sym->ts.deferred
+		      && n->sym->attr.omp_declare_target
+		      && (always_modifier || n->sym->attr.pointer)
+		      && op != EXEC_OMP_TARGET_EXIT_DATA
+		      && n->u.map_op != OMP_MAP_DELETE
+		      && n->u.map_op != OMP_MAP_RELEASE)
+		    {
+		      gcc_assert (n->sym->ts.u.cl->backend_decl);
+		      node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (node5, GOMP_MAP_ALWAYS_TO);
+		      OMP_CLAUSE_DECL (node5) = n->sym->ts.u.cl->backend_decl;
+		      OMP_CLAUSE_SIZE (node5)
+			= TYPE_SIZE_UNIT (gfc_charlen_type_node);
+		    }
+
 		  tree present = gfc_omp_check_optional_argument (decl, true);
 		  if (openacc && n->sym->ts.type == BT_CLASS)
 		    {
@@ -3145,13 +3196,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      OMP_CLAUSE_SIZE (node2) = size_int (0);
 		      goto finalize_map_clause;
 		    }
-		  else if (POINTER_TYPE_P (TREE_TYPE (decl))
+		  else if (POINTER_TYPE_P (type)
 			   && (gfc_omp_privatize_by_reference (decl)
 			       || GFC_DECL_GET_SCALAR_POINTER (decl)
 			       || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
 			       || GFC_DECL_CRAY_POINTEE (decl)
-			       || GFC_DESCRIPTOR_TYPE_P
-					     (TREE_TYPE (TREE_TYPE (decl)))
+			       || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))
 			       || (n->sym->ts.type == BT_DERIVED
 				   && (n->sym->ts.u.derived->ts.f90_type
 				       != BT_VOID))))
@@ -3164,7 +3214,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			 to avoid accessing undefined variables.  We cannot use
 			 a temporary variable here as otherwise the replacement
 			 of the variables in omp-low.cc will not work.  */
-		      if (present && GFC_ARRAY_TYPE_P (TREE_TYPE (decl)))
+		      if (present && GFC_ARRAY_TYPE_P (type))
 			{
 			  tree tmp = fold_build2_loc (input_location,
 						      MODIFY_EXPR,
@@ -3181,22 +3231,51 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 							     cond, tmp,
 							     NULL_TREE));
 			}
-		      node4 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
-		      OMP_CLAUSE_DECL (node4) = decl;
-		      OMP_CLAUSE_SIZE (node4) = size_int (0);
+		      /* For descriptor types, the unmapping happens below.  */
+		      if (op != EXEC_OMP_TARGET_EXIT_DATA
+			  || !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+			{
+			  enum gomp_map_kind gmk = GOMP_MAP_POINTER;
+			  if (op == EXEC_OMP_TARGET_EXIT_DATA
+			      && n->u.map_op == OMP_MAP_DELETE)
+			    gmk = GOMP_MAP_DELETE;
+			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+			    gmk = GOMP_MAP_RELEASE;
+			  tree size;
+			  if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+			    size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+			  else
+			    size = size_int (0);
+			  node4 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node4, gmk);
+			  OMP_CLAUSE_DECL (node4) = decl;
+			  OMP_CLAUSE_SIZE (node4) = size;
+			}
 		      decl = build_fold_indirect_ref (decl);
 		      if ((TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE
 			   || gfc_omp_is_optional_argument (orig_decl))
 			  && (GFC_DECL_GET_SCALAR_POINTER (orig_decl)
 			      || GFC_DECL_GET_SCALAR_ALLOCATABLE (orig_decl)))
 			{
+			  enum gomp_map_kind gmk;
+			  if (op == EXEC_OMP_TARGET_EXIT_DATA
+			      && n->u.map_op == OMP_MAP_DELETE)
+			    gmk = GOMP_MAP_DELETE;
+			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+			    gmk = GOMP_MAP_RELEASE;
+			  else
+			    gmk = GOMP_MAP_POINTER;
+			  tree size;
+			  if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+			    size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+			  else
+			    size = size_int (0);
 			  node3 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+			  OMP_CLAUSE_SET_MAP_KIND (node3, gmk);
 			  OMP_CLAUSE_DECL (node3) = decl;
-			  OMP_CLAUSE_SIZE (node3) = size_int (0);
+			  OMP_CLAUSE_SIZE (node3) = size;
 			  decl = build_fold_indirect_ref (decl);
 			}
 		    }
@@ -3210,56 +3289,70 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
 		      ptr = build_fold_indirect_ref (ptr);
 		      OMP_CLAUSE_DECL (node) = ptr;
-		      node2 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
 		      OMP_CLAUSE_DECL (node2) = decl;
 		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      node3 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      if (present)
-			{
-			  ptr = gfc_conv_descriptor_data_get (decl);
-			  ptr = gfc_build_addr_expr (NULL, ptr);
-			  ptr = gfc_build_cond_assign_expr (block, present, ptr,
-							    null_pointer_node);
-			  ptr = build_fold_indirect_ref (ptr);
-			  OMP_CLAUSE_DECL (node3) = ptr;
-			}
+		      if (n->u.map_op == OMP_MAP_DELETE)
+			map_kind = GOMP_MAP_DELETE;
+		      else if (op == EXEC_OMP_TARGET_EXIT_DATA
+			       || n->u.map_op == OMP_MAP_RELEASE)
+			map_kind = GOMP_MAP_RELEASE;
 		      else
-			OMP_CLAUSE_DECL (node3)
-			  = gfc_conv_descriptor_data_get (decl);
-		      OMP_CLAUSE_SIZE (node3) = size_int (0);
-		      if (n->u.map_op == OMP_MAP_ATTACH)
-			{
-			  /* Standalone attach clauses used with arrays with
-			     descriptors must copy the descriptor to the target,
-			     else they won't have anything to perform the
-			     attachment onto (see OpenACC 2.6, "2.6.3. Data
-			     Structures with Pointers").  */
-			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
-			  /* We don't want to map PTR at all in this case, so
-			     delete its node and shuffle the others down.  */
-			  node = node2;
-			  node2 = node3;
-			  node3 = NULL;
-			  goto finalize_map_clause;
-			}
-		      else if (n->u.map_op == OMP_MAP_DETACH)
+			map_kind = GOMP_MAP_TO_PSET;
+		      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+
+		      if (op != EXEC_OMP_TARGET_EXIT_DATA
+			  && n->u.map_op != OMP_MAP_DELETE
+			  && n->u.map_op != OMP_MAP_RELEASE)
 			{
-			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
-			  /* Similarly to above, we don't want to unmap PTR
-			     here.  */
-			  node = node2;
-			  node2 = node3;
-			  node3 = NULL;
-			  goto finalize_map_clause;
+			  node3 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  if (present)
+			    {
+			      ptr = gfc_conv_descriptor_data_get (decl);
+			      ptr = gfc_build_addr_expr (NULL, ptr);
+			      ptr = gfc_build_cond_assign_expr (
+				      block, present, ptr, null_pointer_node);
+			      ptr = build_fold_indirect_ref (ptr);
+			      OMP_CLAUSE_DECL (node3) = ptr;
+			    }
+			  else
+			    OMP_CLAUSE_DECL (node3)
+			      = gfc_conv_descriptor_data_get (decl);
+			  OMP_CLAUSE_SIZE (node3) = size_int (0);
+
+			  if (n->u.map_op == OMP_MAP_ATTACH)
+			    {
+			      /* Standalone attach clauses used with arrays with
+				 descriptors must copy the descriptor to the
+				 target, else they won't have anything to
+				 perform the attachment onto (see OpenACC 2.6,
+				 "2.6.3. Data Structures with Pointers").  */
+			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
+			      /* We don't want to map PTR at all in this case,
+				 so delete its node and shuffle the others
+				 down.  */
+			      node = node2;
+			      node2 = node3;
+			      node3 = NULL;
+			      goto finalize_map_clause;
+			    }
+			  else if (n->u.map_op == OMP_MAP_DETACH)
+			    {
+			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+			      /* Similarly to above, we don't want to unmap PTR
+				 here.  */
+			      node = node2;
+			      node2 = node3;
+			      node3 = NULL;
+			      goto finalize_map_clause;
+			    }
+			  else
+			    OMP_CLAUSE_SET_MAP_KIND (node3,
+						     always_modifier
+						     ? GOMP_MAP_ALWAYS_POINTER
+						     : GOMP_MAP_POINTER);
 			}
-		      else
-			OMP_CLAUSE_SET_MAP_KIND (node3,
-						 always_modifier
-						 ? GOMP_MAP_ALWAYS_POINTER
-						 : GOMP_MAP_POINTER);
 
 		      /* We have to check for n->sym->attr.dimension because
 			 of scalar coarrays.  */
@@ -3275,6 +3368,23 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  tem
 			    = gfc_full_array_size (&cond_block, decl,
 						   GFC_TYPE_ARRAY_RANK (type));
+			  tree elemsz;
+			  if (n->sym->ts.type == BT_CHARACTER
+			      && n->sym->ts.deferred)
+			    {
+			      tree len = n->sym->ts.u.cl->backend_decl;
+			      len = fold_convert (size_type_node, len);
+			      elemsz = gfc_get_char_type (n->sym->ts.kind);
+			      elemsz = TYPE_SIZE_UNIT (elemsz);
+			      elemsz = fold_build2 (MULT_EXPR, size_type_node,
+						    len, elemsz);
+			    }
+			  else
+			    elemsz
+			      = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+			  elemsz = fold_convert (gfc_array_index_type, elemsz);
+			  tem = fold_build2 (MULT_EXPR, gfc_array_index_type,
+					     tem, elemsz);
 			  gfc_add_modify (&cond_block, size, tem);
 			  then_b = gfc_finish_block (&cond_block);
 			  gfc_init_block (&cond_block);
@@ -3305,6 +3415,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  gfc_init_block (&cond_block);
 			  tree size = gfc_full_array_size (&cond_block, decl,
 					GFC_TYPE_ARRAY_RANK (type));
+			  tree elemsz
+			    = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+			  elemsz = fold_convert (gfc_array_index_type, elemsz);
+			  size = fold_build2 (MULT_EXPR, gfc_array_index_type,
+					      size, elemsz);
+			  size = gfc_evaluate_now (size, &cond_block);
 			  if (present)
 			    {
 			      tree var = gfc_create_var (gfc_array_index_type,
@@ -3323,15 +3439,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      OMP_CLAUSE_SIZE (node) = size;
 			    }
 			}
-		      if (n->sym->attr.dimension)
-			{
-			  tree elemsz
-			    = TYPE_SIZE_UNIT (gfc_get_element_type (type));
-			  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);
-			}
 		    }
 		  else if (present
 			   && TREE_CODE (decl) == INDIRECT_REF
@@ -3347,6 +3454,37 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		    }
 		  else
 		    OMP_CLAUSE_DECL (node) = decl;
+
+		  if (!n->sym->attr.dimension
+		      && n->sym->ts.type == BT_CHARACTER
+		      && n->sym->ts.deferred)
+		    {
+		      if (!DECL_P (decl))
+			{
+			  gcc_assert (TREE_CODE (decl) == INDIRECT_REF);
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		      tree cond = fold_build2_loc (input_location, NE_EXPR,
+						   boolean_type_node,
+						   decl, null_pointer_node);
+		      if (present)
+			cond = fold_build2_loc (input_location,
+						TRUTH_ANDIF_EXPR,
+						boolean_type_node,
+						present, cond);
+		      tree len = n->sym->ts.u.cl->backend_decl;
+		      len = fold_convert (size_type_node, len);
+		      tree size = gfc_get_char_type (n->sym->ts.kind);
+		      size = TYPE_SIZE_UNIT (size);
+		      size = fold_build2 (MULT_EXPR, size_type_node, len, size);
+		      size = build3_loc (input_location,
+							 COND_EXPR,
+							 size_type_node,
+							 cond, size,
+							 size_zero_node);
+		      size = gfc_evaluate_now (size, block);
+		      OMP_CLAUSE_SIZE (node) = size;
+		    }
 		}
 	      else if (n->expr
 		       && n->expr->expr_type == EXPR_VARIABLE
@@ -3363,7 +3501,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      && !(POINTER_TYPE_P (type)
 			   && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))))
 		    k = GOMP_MAP_FIRSTPRIVATE_POINTER;
-		  gfc_trans_omp_array_section (block, n, decl, element, k,
+		  gfc_trans_omp_array_section (block, op, n, decl, element, k,
 					       node, node2, node3, node4);
 		}
 	      else if (n->expr
@@ -3424,9 +3562,15 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 					   fold_convert (size_type_node,
 					       se.string_length),
 					   TYPE_SIZE_UNIT (tmp));
+			  if (n->u.map_op == OMP_MAP_DELETE)
+			    kind = GOMP_MAP_DELETE;
+			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+			    kind = GOMP_MAP_RELEASE;
+			  else
+			    kind = GOMP_MAP_TO;
 			  node3 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO);
+			  OMP_CLAUSE_SET_MAP_KIND (node3, kind);
 			  OMP_CLAUSE_DECL (node3) = se.string_length;
 			  OMP_CLAUSE_SIZE (node3)
 			    = TYPE_SIZE_UNIT (gfc_charlen_type_node);
@@ -3551,11 +3695,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    = gfc_full_array_size (block, inner, rank);
 			  tree elemsz
 			    = TYPE_SIZE_UNIT (gfc_get_element_type (type));
-			  if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
-			    map_kind = GOMP_MAP_TO;
+			  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)
-			    map_kind = OMP_CLAUSE_MAP_KIND (node);
+			    ;
+			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+			    map_kind = GOMP_MAP_RELEASE;
 			  else
 			    map_kind = GOMP_MAP_ALLOC;
 			  if (!openacc
@@ -3596,6 +3746,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      node2 = node;
 			      node = desc_node;  /* Put first.  */
 			    }
+			  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,
@@ -3626,7 +3778,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      bool element = lastref->u.ar.type == AR_ELEMENT;
 		      gomp_map_kind kind = (openacc ? GOMP_MAP_ATTACH_DETACH
 						    : GOMP_MAP_ALWAYS_POINTER);
-		      gfc_trans_omp_array_section (block, n, inner, element,
+		      gfc_trans_omp_array_section (block, op, n, inner, element,
 						   kind, node, node2, node3,
 						   node4);
 		    }
@@ -3645,6 +3797,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
 	      if (node4)
 		omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+	      if (node5)
+		omp_clauses = gfc_trans_add_clause (node5, omp_clauses);
 	    }
 	  break;
 	case OMP_LIST_TO:
@@ -7512,7 +7666,7 @@ gfc_trans_omp_target_exit_data (gfc_code *code)
 
   gfc_start_block (&block);
   omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-				       code->loc);
+				       code->loc, false, false, code->op);
   stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node,
 		     omp_clauses);
   gfc_add_expr_to_block (&block, stmt);
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 3740a8979af..7afa4c31a2b 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10741,7 +10741,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 {
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
-  tree *prev_list_p = NULL, *orig_list_p = list_p;
+  tree *orig_list_p = list_p;
   int handled_depend_iterators = -1;
   int nowait = -1;
 
@@ -11168,31 +11168,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    default:
 	      break;
 	    }
-	  /* For Fortran, not only the pointer to the data is mapped but also
-	     the address of the pointer, the array descriptor etc.; for
-	     'exit data' - and in particular for 'delete:' - having an 'alloc:'
-	     does not make sense.  Likewise, for 'update' only transferring the
-	     data itself is needed as the rest has been handled in previous
-	     directives.  However, for 'exit data', the array descriptor needs
-	     to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
-
-	     NOTE: Generally, it is not safe to perform "enter data" operations
-	     on arrays where the data *or the descriptor* may go out of scope
-	     before a corresponding "exit data" operation -- and such a
-	     descriptor may be synthesized temporarily, e.g. to pass an
-	     explicit-shape array to a function expecting an assumed-shape
-	     argument.  Performing "enter data" inside the called function
-	     would thus be problematic.  */
-	  if (code == OMP_TARGET_EXIT_DATA
-	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
-	    OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
-					== GOMP_MAP_DELETE
-					? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
-	  else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
-		   && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
-		       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
-	    remove = true;
-
 	  if (remove)
 	    break;
 	  if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
@@ -11452,21 +11427,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  remove = true;
 		  break;
 		}
-
-	      if (!remove
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
-		  && OMP_CLAUSE_CHAIN (c)
-		  && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
-		  && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-		       == GOMP_MAP_ALWAYS_POINTER)
-		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			  == GOMP_MAP_ATTACH_DETACH)
-		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			  == GOMP_MAP_TO_PSET)))
-		prev_list_p = list_p;
-
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 1e5bf0ba1e6..23f0ffc627e 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\\(to: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\\(to:del_f_p \\\[pointer set, 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 \\\[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" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-10.f90 b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
new file mode 100644
index 00000000000..c12bf25ad19
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
@@ -0,0 +1,69 @@
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! If enter data adds a (GOMP_MAP_)POINTER attachment, exit data needs to remove
+! it again. If not there can be all kind of issues, in particular when
+! stack memory was mapped, reused later and mapped again.
+
+subroutine test_aa (aa2, aa3)
+  integer(kind=4), allocatable :: aa1, aa2, aa3
+  optional :: aa3
+  !$omp target enter data map(aa1)
+  !$omp target exit data map(aa1)
+  !$omp target enter data map(aa2)
+  !$omp target exit data map(aa2)
+  !$omp target enter data map(aa3)
+  !$omp target exit data map(aa3)
+end
+
+subroutine test_pp (pp2, pp3)
+  integer(kind=4), allocatable :: pp1, pp2, pp3
+  optional :: pp3
+  !$omp target enter data map(pp1)
+  !$omp target exit data map(pp1)
+  !$omp target enter data map(pp2)
+  !$omp target exit data map(pp2)
+  !$omp target enter data map(pp3)
+  !$omp target exit data map(pp3)
+end
+
+subroutine test_pprelease (rp2, rp3)
+  integer(kind=4), allocatable :: rp1, rp2, rp3
+  optional :: rp3
+  !$omp target enter data map(rp1)
+  !$omp target exit data map(release:rp1)
+  !$omp target enter data map(rp2)
+  !$omp target exit data map(release:rp2)
+  !$omp target enter data map(rp3)
+  !$omp target exit data map(release:rp3)
+end
+
+subroutine test_ppdelete (dp2, dp3)
+  integer(kind=4), allocatable :: dp1, dp2, dp3
+  optional :: dp3
+  !$omp target enter data map(dp1)
+  !$omp target exit data map(delete:dp1)
+  !$omp target enter data map(dp2)
+  !$omp target exit data map(delete:dp2)
+  !$omp target enter data map(dp3)
+  !$omp target exit data map(delete:dp3)
+end
+
+
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:aa1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa2 \\\[len: .\\\]\\) map\\(release:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa3 \\\[len: .\\\]\\) map\\(release:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:pp1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp2 \\\[len: .\\\]\\) map\\(release:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp3 \\\[len: .\\\]\\) map\\(release:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp1 \\\[len: .\\\]\\) map\\(release:\\*rp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp2 \\\[len: .\\\]\\) map\\(release:\\*rp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:rp3 \\\[len: .\\\]\\) map\\(release:\\*rp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp1 \\\[len: .\\\]\\) map\\(delete:\\*dp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp2 \\\[len: .\\\]\\) map\\(delete:\\*dp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(delete:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:dp3 \\\[len: .\\\]\\) map\\(delete:\\*dp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 b/gcc/testsuite/gfortran.dg/gomp/map-9.f90
index 9e7b811c8af..b770b931bee 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\\(to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { 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\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
 
 program p
    type t
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
index f5d888592b9..cd771b33a25 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
@@ -48,9 +48,11 @@ contains
   end subroutine sub
 end module m
 
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! Check for multiplication: len = arrays_size * 4<bytes>:
+! { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = D\\.\[0-9\]+ \\* 4;" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90 b/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90
index ed57d0072d7..219dc467c46 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90
@@ -15,6 +15,6 @@ integer, allocatable :: one(:), two(:), three(:)
 !$omp target exit data map(from:three)
 end
 
-! { dg-final { scan-tree-dump "omp target exit data map\\(delete:.*\\) map\\(delete:one \\\[len: .*\\\]\\)" "omplower" } }
-! { dg-final { scan-tree-dump "omp target exit data map\\(release:.*\\) map\\(release:two \\\[len: .*\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:one \\\[len: \[0-9\]+\\\]\\) map\\(delete:MEM " "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:two \\\[len: \[0-9\]+\\\]\\) map\\(release:MEM " "omplower" } }
 ! { dg-final { scan-tree-dump "omp target exit data map\\(from:.*\\) map\\(release:three \\\[len: .*\\\]\\)" "omplower" } }
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90
index 5d97566c66c..1b3cdf9e76a 100644
--- a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90
@@ -17,6 +17,6 @@ var%p2 = [46,679,54]
   if (any (var%p1 /= [22,53,28,6,4])) stop 3
   if (any (var%p2 /= [46,679,54])) stop 4
 !$omp end target
-!!$omp target exit data map(from:var%p1, var%p2)
+!$omp target exit data map(from:var%p1, var%p2)
 end
 
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90
new file mode 100644
index 00000000000..6192bf29d9a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90
@@ -0,0 +1,540 @@
+! Check that 'map(alloc:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+  integer :: ic(2:5), ic2
+  character(len=11) :: ccstr(3:4), ccstr2
+  character(len=11,kind=4) :: cc4str(3:7), cc4str2
+  integer, pointer :: pc(:), pc2
+  character(len=:), pointer :: pcstr(:), pcstr2
+  character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+end type t
+
+type(t) :: dt
+
+integer :: ii(5), ii2
+character(len=11) :: clstr(-1:1), clstr2
+character(len=11,kind=4) :: cl4str(0:3), cl4str2
+integer, pointer :: ip(:), ip2
+integer, allocatable :: ia(:), ia2
+character(len=:), pointer :: pstr(:), pstr2
+character(len=:), allocatable :: astr(:), astr2
+character(len=:,kind=4), pointer :: p4str(:), p4str2
+character(len=:,kind=4), allocatable :: a4str(:), a4str2
+
+allocate(dt%pc(5), dt%pc2)
+allocate(character(len=2) :: dt%pcstr(2))
+allocate(character(len=4) :: dt%pcstr2)
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+allocate(ip(5), ip2, ia(8), ia2)
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=4) :: pstr2)
+allocate(character(len=6) :: astr(3:5))
+allocate(character(len=8) :: astr2)
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=5,kind=4) :: p4str2)
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+allocate(character(len=9,kind=4) :: a4str2)
+
+
+! integer :: ic(2:5), ic2
+
+!$omp target enter data map(alloc: dt%ic)
+!$omp target map(alloc: dt%ic)
+  if (size(dt%ic) /= 4) error stop
+  if (lbound(dt%ic, 1) /= 2) error stop
+  if (ubound(dt%ic, 1) /= 5) error stop
+  dt%ic = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic)
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: dt%ic2)
+!$omp target map(alloc: dt%ic2)
+  dt%ic2 = 42
+!$omp end target
+!$omp target exit data map(from: dt%ic2)
+if (dt%ic2 /= 42) error stop
+
+
+! character(len=11) :: ccstr(3:4), ccstr2
+
+!$omp target enter data map(alloc: dt%ccstr)
+!$omp target map(alloc: dt%ccstr)
+  if (len(dt%ccstr) /= 11) error stop
+  if (size(dt%ccstr) /= 2) error stop
+  if (lbound(dt%ccstr, 1) /= 3) error stop
+  if (ubound(dt%ccstr, 1) /= 4) error stop
+  dt%ccstr = ["12345678901", "abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr)
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+!$omp target enter data map(alloc: dt%ccstr2)
+!$omp target map(alloc: dt%ccstr2)
+  if (len(dt%ccstr2) /= 11) error stop
+  dt%ccstr2 = "ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%ccstr2)
+if (len(dt%ccstr2) /= 11) error stop
+if (dt%ccstr2 /= "ABCDEFGHIJK") error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7), cc4str2
+
+! Value check fails
+!$omp target enter data map(alloc: dt%cc4str)
+!$omp target map(alloc: dt%cc4str)
+  if (len(dt%cc4str) /= 11) error stop
+  if (size(dt%cc4str) /= 5) error stop
+  if (lbound(dt%cc4str, 1) /= 3) error stop
+  if (ubound(dt%cc4str, 1) /= 7) error stop
+  dt%cc4str = [4_"12345678901", 4_"abcdefghijk", &
+               4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+               4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str)
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+!$omp target enter data map(alloc: dt%cc4str2)
+!$omp target map(alloc: dt%cc4str2)
+  if (len(dt%cc4str2) /= 11) error stop
+  dt%cc4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%cc4str2)
+if (len(dt%cc4str2) /= 11) error stop
+if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! integer, pointer :: pc(:), pc2
+! allocate(dt%pc(5), dt%pc2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pc)
+!$omp target map(alloc: dt%pc)
+  if (.not. associated(dt%pc)) error stop
+  if (size(dt%pc) /= 5) error stop
+  if (lbound(dt%pc, 1) /= 1) error stop
+  if (ubound(dt%pc, 1) /= 5) error stop
+  dt%pc = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc)
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: dt%pc2)
+!$omp target map(alloc: dt%pc2)
+  if (.not. associated(dt%pc2)) error stop
+  dt%pc2 = 99
+!$omp end target
+!$omp target exit data map(from: dt%pc2)
+if (dt%pc2 /= 99) error stop
+if (.not. associated(dt%pc2)) error stop
+
+
+! character(len=:), pointer :: pcstr(:), pcstr2
+! allocate(character(len=2) :: dt%pcstr(2))
+! allocate(character(len=4) :: dt%pcstr2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pcstr)
+!$omp target map(alloc: dt%pcstr)
+  if (.not. associated(dt%pcstr)) error stop
+  if (len(dt%pcstr) /= 2) error stop
+  if (size(dt%pcstr) /= 2) error stop
+  if (lbound(dt%pcstr, 1) /= 1) error stop
+  if (ubound(dt%pcstr, 1) /= 2) error stop
+  dt%pcstr = ["01", "jk"]
+!$omp end target
+!$omp target exit data map(from: dt%pcstr)
+if (.not. associated(dt%pcstr)) error stop
+if (len(dt%pcstr) /= 2) error stop
+if (size(dt%pcstr) /= 2) error stop
+if (lbound(dt%pcstr, 1) /= 1) error stop
+if (ubound(dt%pcstr, 1) /= 2) error stop
+if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(alloc: dt%pcstr2)
+!$omp target map(alloc: dt%pcstr2)
+  if (.not. associated(dt%pcstr2)) error stop
+  if (len(dt%pcstr2) /= 4) error stop
+  dt%pcstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: dt%pcstr2)
+if (.not. associated(dt%pcstr2)) error stop
+if (len(dt%pcstr2) /= 4) error stop
+if (dt%pcstr2 /= "HIJK") error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+! allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+! structure element when other mapped elements from the same structure weren't mapped together with it
+!$omp target enter data map(alloc: dt%pc4str)
+!$omp target map(alloc: dt%pc4str)
+  if (.not. associated(dt%pc4str)) error stop
+  if (len(dt%pc4str) /= 3) error stop
+  if (size(dt%pc4str) /= 2) error stop
+  if (lbound(dt%pc4str, 1) /= 2) error stop
+  if (ubound(dt%pc4str, 1) /= 3) error stop
+  dt%pc4str = [4_"456", 4_"tzu"]
+!$omp end target
+!$omp target exit data map(from: dt%pc4str)
+if (.not. associated(dt%pc4str)) error stop
+if (len(dt%pc4str) /= 3) error stop
+if (size(dt%pc4str) /= 2) error stop
+if (lbound(dt%pc4str, 1) /= 2) error stop
+if (ubound(dt%pc4str, 1) /= 3) error stop
+if (dt%pc4str(2) /= 4_"456") error stop
+if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(alloc: dt%pc4str2)
+!$omp target map(alloc: dt%pc4str2)
+  if (.not. associated(dt%pc4str2)) error stop
+  if (len(dt%pc4str2) /= 5) error stop
+  dt%pc4str2 = 4_"98765"
+!$omp end target
+!$omp target exit data map(from: dt%pc4str2)
+if (.not. associated(dt%pc4str2)) error stop
+if (len(dt%pc4str2) /= 5) error stop
+if (dt%pc4str2 /= 4_"98765") error stop
+
+
+! integer :: ii(5), ii2
+
+!$omp target enter data map(alloc: ii)
+!$omp target map(alloc: ii)
+  if (size(ii) /= 5) error stop
+  if (lbound(ii, 1) /= 1) error stop
+  if (ubound(ii, 1) /= 5) error stop
+  ii = [-1, -2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii)
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+!$omp target enter data map(alloc: ii2)
+!$omp target map(alloc: ii2)
+  ii2 = -410
+!$omp end target
+!$omp target exit data map(from: ii2)
+if (ii2 /= -410) error stop
+
+
+! character(len=11) :: clstr(-1:1), clstr2
+
+!$omp target enter data map(alloc: clstr)
+!$omp target map(alloc: clstr)
+  if (len(clstr) /= 11) error stop
+  if (size(clstr) /= 3) error stop
+  if (lbound(clstr, 1) /= -1) error stop
+  if (ubound(clstr, 1) /= 1) error stop
+  clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr)
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+!$omp target enter data map(alloc: clstr2)
+!$omp target map(alloc: clstr2)
+  if (len(clstr2) /= 11) error stop
+  clstr2 = "ABCDEFghijk"
+!$omp end target
+!$omp target exit data map(from: clstr2)
+if (len(clstr2) /= 11) error stop
+if (clstr2 /= "ABCDEFghijk") error stop
+
+
+! character(len=11,kind=4) :: cl4str(0:3), cl4str2
+
+!$omp target enter data map(alloc: cl4str)
+!$omp target map(alloc: cl4str)
+  if (len(cl4str) /= 11) error stop
+  if (size(cl4str) /= 4) error stop
+  if (lbound(cl4str, 1) /= 0) error stop
+  if (ubound(cl4str, 1) /= 3) error stop
+  cl4str = [4_"12345678901", 4_"abcdefghijk", &
+            4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str)
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+!$omp target enter data map(alloc: cl4str2)
+!$omp target map(alloc: cl4str2)
+  if (len(cl4str2) /= 11) error stop
+  cl4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: cl4str2)
+if (len(cl4str2) /= 11) error stop
+if (cl4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(alloc: ip)
+!$omp target map(alloc: ip)
+  if (.not. associated(ip)) error stop
+  if (size(ip) /= 5) error stop
+  if (lbound(ip, 1) /= 1) error stop
+  if (ubound(ip, 1) /= 5) error stop
+  ip = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip)
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: ip2)
+!$omp target map(alloc: ip2)
+  if (.not. associated(ip2)) error stop
+  ip2 = 99
+!$omp end target
+!$omp target exit data map(from: ip2)
+if (ip2 /= 99) error stop
+if (.not. associated(ip2)) error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(alloc: ia)
+!$omp target map(alloc: ia)
+  if (.not. allocated(ia)) error stop
+  if (size(ia) /= 8) error stop
+  if (lbound(ia, 1) /= 1) error stop
+  if (ubound(ia, 1) /= 8) error stop
+  ia = [1,2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia)
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+!$omp target enter data map(alloc: ia2)
+!$omp target map(alloc: ia2)
+  if (.not. allocated(ia2)) error stop
+  ia2 = 102
+!$omp end target
+!$omp target exit data map(from: ia2)
+if (ia2 /= 102) error stop
+if (.not. allocated(ia2)) error stop
+
+
+! character(len=:), pointer :: pstr(:), pstr2
+! allocate(character(len=2) :: pstr(-2:0))
+! allocate(character(len=4) :: pstr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: pstr)
+!$omp target map(alloc: pstr)
+  if (.not. associated(pstr)) error stop
+  if (len(pstr) /= 2) error stop
+  if (size(pstr) /= 3) error stop
+  if (lbound(pstr, 1) /= -2) error stop
+  if (ubound(pstr, 1) /= 0) error stop
+  pstr = ["01", "jk", "aq"]
+!$omp end target
+!$omp target exit data map(from: pstr)
+if (.not. associated(pstr)) error stop
+if (len(pstr) /= 2) error stop
+if (size(pstr) /= 3) error stop
+if (lbound(pstr, 1) /= -2) error stop
+if (ubound(pstr, 1) /= 0) error stop
+if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+!$omp target enter data map(alloc: pstr2)
+!$omp target map(alloc: pstr2)
+  if (.not. associated(pstr2)) error stop
+  if (len(pstr2) /= 4) error stop
+  pstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: pstr2)
+if (.not. associated(pstr2)) error stop
+if (len(pstr2) /= 4) error stop
+if (pstr2 /= "HIJK") error stop
+
+
+! character(len=:), allocatable :: astr(:), astr2
+! allocate(character(len=6) :: astr(3:5))
+! allocate(character(len=8) :: astr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: astr)
+!$omp target map(alloc: astr)
+  if (.not. allocated(astr)) error stop
+  if (len(astr) /= 6) error stop
+  if (size(astr) /= 3) error stop
+  if (lbound(astr, 1) /= 3) error stop
+  if (ubound(astr, 1) /= 5) error stop
+  astr = ["01db45", "jk$D%S", "zutg47"]
+!$omp end target
+!$omp target exit data map(from: astr)
+if (.not. allocated(astr)) error stop
+if (len(astr) /= 6) error stop
+if (size(astr) /= 3) error stop
+if (lbound(astr, 1) /= 3) error stop
+if (ubound(astr, 1) /= 5) error stop
+if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: astr2)
+!$omp target map(alloc: astr2)
+  if (.not. allocated(astr2)) error stop
+  if (len(astr2) /= 8) error stop
+  astr2 = "HIJKhijk"
+!$omp end target
+!$omp target exit data map(from: astr2)
+if (.not. allocated(astr2)) error stop
+if (len(astr2) /= 8) error stop
+if (astr2 /= "HIJKhijk") error stop
+
+
+! character(len=:,kind=4), pointer :: p4str(:), p4str2
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+! allocate(character(len=5,kind=4) :: p4str2)
+
+! FAILS with value check
+
+!$omp target enter data map(alloc: p4str)
+!$omp target map(alloc: p4str)
+  if (.not. associated(p4str)) error stop
+  if (len(p4str) /= 3) error stop
+  if (size(p4str) /= 3) error stop
+  if (lbound(p4str, 1) /= 2) error stop
+  if (ubound(p4str, 1) /= 4) error stop
+  p4str(:) = [4_"f85", 4_"8af", 4_"A%F"]
+!$omp end target
+!$omp target exit data map(from: p4str)
+if (.not. associated(p4str)) error stop
+if (len(p4str) /= 3) error stop
+if (size(p4str) /= 3) error stop
+if (lbound(p4str, 1) /= 2) error stop
+if (ubound(p4str, 1) /= 4) error stop
+if (p4str(2)  /= 4_"f85") error stop
+if (p4str(3)  /= 4_"8af") error stop
+if (p4str(4)  /= 4_"A%F") error stop
+
+!$omp target enter data map(alloc: p4str2)
+!$omp target map(alloc: p4str2)
+  if (.not. associated(p4str2)) error stop
+  if (len(p4str2) /= 5) error stop
+  p4str2 = 4_"9875a"
+!$omp end target
+!$omp target exit data map(from: p4str2)
+if (.not. associated(p4str2)) error stop
+if (len(p4str2) /= 5) error stop
+if (p4str2 /= 4_"9875a") error stop
+
+
+! character(len=:,kind=4), allocatable :: a4str(:), a4str2
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+! allocate(character(len=9,kind=4) :: a4str2)
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+!$omp target enter data map(alloc: a4str)
+!$omp target map(alloc: a4str)
+  if (.not. allocated(a4str)) error stop
+  if (len(a4str) /= 7) error stop
+  if (size(a4str) /= 6) error stop
+  if (lbound(a4str, 1) /= -2) error stop
+  if (ubound(a4str, 1) /= 3) error stop
+  ! See PR fortran/107508 why '(:)' is required
+  a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!$omp end target
+!$omp target exit data map(from: a4str)
+if (.not. allocated(a4str)) error stop
+if (len(a4str) /= 7) error stop
+if (size(a4str) /= 6) error stop
+if (lbound(a4str, 1) /= -2) error stop
+if (ubound(a4str, 1) /= 3) error stop
+if (a4str(-2) /= 4_"sf456aq") error stop
+if (a4str(-1) /= 4_"3dtzu24") error stop
+if (a4str(0)  /= 4_"_4fh7sm") error stop
+if (a4str(1)  /= 4_"=ff85s7") error stop
+if (a4str(2)  /= 4_"j=8af4d") error stop
+if (a4str(3)  /= 4_".,A%Fsz") error stop
+
+!$omp target enter data map(alloc: a4str2)
+!$omp target map(alloc: a4str2)
+  if (.not. allocated(a4str2)) error stop
+  if (len(a4str2) /= 9) error stop
+  a4str2 = 4_"98765a23d"
+!$omp end target
+!$omp target exit data map(from: a4str2)
+if (.not. allocated(a4str2)) error stop
+if (len(a4str2) /= 9) error stop
+if (a4str2 /= 4_"98765a23d") error stop
+
+
+deallocate(dt%pc, dt%pc2)
+deallocate(dt%pcstr)
+deallocate(dt%pcstr2)
+
+deallocate(dt%pc4str)
+deallocate(dt%pc4str2)
+
+deallocate(ip, ip2, ia, ia2)
+deallocate(pstr)
+deallocate(pstr2)
+deallocate(astr)
+deallocate(astr2)
+
+deallocate(p4str)
+deallocate(p4str2)
+deallocate(a4str)
+deallocate(a4str2)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90
new file mode 100644
index 00000000000..cf759346a84
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90
@@ -0,0 +1,540 @@
+! Check that 'map((to)from:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+  integer :: ic(2:5), ic2
+  character(len=11) :: ccstr(3:4), ccstr2
+  character(len=11,kind=4) :: cc4str(3:7), cc4str2
+  integer, pointer :: pc(:), pc2
+  character(len=:), pointer :: pcstr(:), pcstr2
+  character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+end type t
+
+type(t) :: dt
+
+integer :: ii(5), ii2
+character(len=11) :: clstr(-1:1), clstr2
+character(len=11,kind=4) :: cl4str(0:3), cl4str2
+integer, pointer :: ip(:), ip2
+integer, allocatable :: ia(:), ia2
+character(len=:), pointer :: pstr(:), pstr2
+character(len=:), allocatable :: astr(:), astr2
+character(len=:,kind=4), pointer :: p4str(:), p4str2
+character(len=:,kind=4), allocatable :: a4str(:), a4str2
+
+allocate(dt%pc(5), dt%pc2)
+allocate(character(len=2) :: dt%pcstr(2))
+allocate(character(len=4) :: dt%pcstr2)
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+allocate(ip(5), ip2, ia(8), ia2)
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=4) :: pstr2)
+allocate(character(len=6) :: astr(3:5))
+allocate(character(len=8) :: astr2)
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=5,kind=4) :: p4str2)
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+allocate(character(len=9,kind=4) :: a4str2)
+
+
+! integer :: ic(2:5), ic2
+
+!$omp target enter data map(tofrom: dt%ic)
+!$omp target map(from: dt%ic)
+  if (size(dt%ic) /= 4) error stop
+  if (lbound(dt%ic, 1) /= 2) error stop
+  if (ubound(dt%ic, 1) /= 5) error stop
+  dt%ic = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic)
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: dt%ic2)
+!$omp target map(from: dt%ic2)
+  dt%ic2 = 42
+!$omp end target
+!$omp target exit data map(from: dt%ic2)
+if (dt%ic2 /= 42) error stop
+
+
+! character(len=11) :: ccstr(3:4), ccstr2
+
+!$omp target enter data map(tofrom: dt%ccstr)
+!$omp target map(from: dt%ccstr)
+  if (len(dt%ccstr) /= 11) error stop
+  if (size(dt%ccstr) /= 2) error stop
+  if (lbound(dt%ccstr, 1) /= 3) error stop
+  if (ubound(dt%ccstr, 1) /= 4) error stop
+  dt%ccstr = ["12345678901", "abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr)
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+!$omp target enter data map(tofrom: dt%ccstr2)
+!$omp target map(from: dt%ccstr2)
+  if (len(dt%ccstr2) /= 11) error stop
+  dt%ccstr2 = "ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%ccstr2)
+if (len(dt%ccstr2) /= 11) error stop
+if (dt%ccstr2 /= "ABCDEFGHIJK") error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7), cc4str2
+
+! Value check fails
+!$omp target enter data map(tofrom: dt%cc4str)
+!$omp target map(from: dt%cc4str)
+  if (len(dt%cc4str) /= 11) error stop
+  if (size(dt%cc4str) /= 5) error stop
+  if (lbound(dt%cc4str, 1) /= 3) error stop
+  if (ubound(dt%cc4str, 1) /= 7) error stop
+  dt%cc4str = [4_"12345678901", 4_"abcdefghijk", &
+               4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+               4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str)
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+!$omp target enter data map(tofrom: dt%cc4str2)
+!$omp target map(from: dt%cc4str2)
+  if (len(dt%cc4str2) /= 11) error stop
+  dt%cc4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%cc4str2)
+if (len(dt%cc4str2) /= 11) error stop
+if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! integer, pointer :: pc(:), pc2
+! allocate(dt%pc(5), dt%pc2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pc)
+!$omp target map(from: dt%pc)
+  if (.not. associated(dt%pc)) error stop
+  if (size(dt%pc) /= 5) error stop
+  if (lbound(dt%pc, 1) /= 1) error stop
+  if (ubound(dt%pc, 1) /= 5) error stop
+  dt%pc = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc)
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: dt%pc2)
+!$omp target map(from: dt%pc2)
+  if (.not. associated(dt%pc2)) error stop
+  dt%pc2 = 99
+!$omp end target
+!$omp target exit data map(from: dt%pc2)
+if (dt%pc2 /= 99) error stop
+if (.not. associated(dt%pc2)) error stop
+
+
+! character(len=:), pointer :: pcstr(:), pcstr2
+! allocate(character(len=2) :: dt%pcstr(2))
+! allocate(character(len=4) :: dt%pcstr2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pcstr)
+!$omp target map(from: dt%pcstr)
+  if (.not. associated(dt%pcstr)) error stop
+  if (len(dt%pcstr) /= 2) error stop
+  if (size(dt%pcstr) /= 2) error stop
+  if (lbound(dt%pcstr, 1) /= 1) error stop
+  if (ubound(dt%pcstr, 1) /= 2) error stop
+  dt%pcstr = ["01", "jk"]
+!$omp end target
+!$omp target exit data map(from: dt%pcstr)
+if (.not. associated(dt%pcstr)) error stop
+if (len(dt%pcstr) /= 2) error stop
+if (size(dt%pcstr) /= 2) error stop
+if (lbound(dt%pcstr, 1) /= 1) error stop
+if (ubound(dt%pcstr, 1) /= 2) error stop
+if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(tofrom: dt%pcstr2)
+!$omp target map(from: dt%pcstr2)
+  if (.not. associated(dt%pcstr2)) error stop
+  if (len(dt%pcstr2) /= 4) error stop
+  dt%pcstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: dt%pcstr2)
+if (.not. associated(dt%pcstr2)) error stop
+if (len(dt%pcstr2) /= 4) error stop
+if (dt%pcstr2 /= "HIJK") error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+! allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pc4str)
+!$omp target map(from: dt%pc4str)
+  if (.not. associated(dt%pc4str)) error stop
+  if (len(dt%pc4str) /= 3) error stop
+  if (size(dt%pc4str) /= 2) error stop
+  if (lbound(dt%pc4str, 1) /= 2) error stop
+  if (ubound(dt%pc4str, 1) /= 3) error stop
+  dt%pc4str = [4_"456", 4_"tzu"]
+!$omp end target
+!$omp target exit data map(from: dt%pc4str)
+if (.not. associated(dt%pc4str)) error stop
+if (len(dt%pc4str) /= 3) error stop
+if (size(dt%pc4str) /= 2) error stop
+if (lbound(dt%pc4str, 1) /= 2) error stop
+if (ubound(dt%pc4str, 1) /= 3) error stop
+if (dt%pc4str(2) /= 4_"456") error stop
+if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(tofrom: dt%pc4str2)
+!$omp target map(from: dt%pc4str2)
+  if (.not. associated(dt%pc4str2)) error stop
+  if (len(dt%pc4str2) /= 5) error stop
+  dt%pc4str2 = 4_"98765"
+!$omp end target
+!$omp target exit data map(from: dt%pc4str2)
+if (.not. associated(dt%pc4str2)) error stop
+if (len(dt%pc4str2) /= 5) error stop
+if (dt%pc4str2 /= 4_"98765") error stop
+
+
+! integer :: ii(5), ii2
+
+!$omp target enter data map(tofrom: ii)
+!$omp target map(from: ii)
+  if (size(ii) /= 5) error stop
+  if (lbound(ii, 1) /= 1) error stop
+  if (ubound(ii, 1) /= 5) error stop
+  ii = [-1, -2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii)
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+!$omp target enter data map(tofrom: ii2)
+!$omp target map(from: ii2)
+  ii2 = -410
+!$omp end target
+!$omp target exit data map(from: ii2)
+if (ii2 /= -410) error stop
+
+
+! character(len=11) :: clstr(-1:1), clstr2
+
+!$omp target enter data map(tofrom: clstr)
+!$omp target map(from: clstr)
+  if (len(clstr) /= 11) error stop
+  if (size(clstr) /= 3) error stop
+  if (lbound(clstr, 1) /= -1) error stop
+  if (ubound(clstr, 1) /= 1) error stop
+  clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr)
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+!$omp target enter data map(tofrom: clstr2)
+!$omp target map(from: clstr2)
+  if (len(clstr2) /= 11) error stop
+  clstr2 = "ABCDEFghijk"
+!$omp end target
+!$omp target exit data map(from: clstr2)
+if (len(clstr2) /= 11) error stop
+if (clstr2 /= "ABCDEFghijk") error stop
+
+
+! character(len=11,kind=4) :: cl4str(0:3), cl4str2
+
+!$omp target enter data map(tofrom: cl4str)
+!$omp target map(from: cl4str)
+  if (len(cl4str) /= 11) error stop
+  if (size(cl4str) /= 4) error stop
+  if (lbound(cl4str, 1) /= 0) error stop
+  if (ubound(cl4str, 1) /= 3) error stop
+  cl4str = [4_"12345678901", 4_"abcdefghijk", &
+            4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str)
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+!$omp target enter data map(tofrom: cl4str2)
+!$omp target map(from: cl4str2)
+  if (len(cl4str2) /= 11) error stop
+  cl4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: cl4str2)
+if (len(cl4str2) /= 11) error stop
+if (cl4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(tofrom: ip)
+!$omp target map(from: ip)
+  if (.not. associated(ip)) error stop
+  if (size(ip) /= 5) error stop
+  if (lbound(ip, 1) /= 1) error stop
+  if (ubound(ip, 1) /= 5) error stop
+  ip = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip)
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: ip2)
+!$omp target map(from: ip2)
+  if (.not. associated(ip2)) error stop
+  ip2 = 99
+!$omp end target
+!$omp target exit data map(from: ip2)
+if (ip2 /= 99) error stop
+if (.not. associated(ip2)) error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(tofrom: ia)
+!$omp target map(from: ia)
+  if (.not. allocated(ia)) error stop
+  if (size(ia) /= 8) error stop
+  if (lbound(ia, 1) /= 1) error stop
+  if (ubound(ia, 1) /= 8) error stop
+  ia = [1,2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia)
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+!$omp target enter data map(tofrom: ia2)
+!$omp target map(from: ia2)
+  if (.not. allocated(ia2)) error stop
+  ia2 = 102
+!$omp end target
+!$omp target exit data map(from: ia2)
+if (ia2 /= 102) error stop
+if (.not. allocated(ia2)) error stop
+
+
+! character(len=:), pointer :: pstr(:), pstr2
+! allocate(character(len=2) :: pstr(-2:0))
+! allocate(character(len=4) :: pstr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: pstr)
+!$omp target map(from: pstr)
+  if (.not. associated(pstr)) error stop
+  if (len(pstr) /= 2) error stop
+  if (size(pstr) /= 3) error stop
+  if (lbound(pstr, 1) /= -2) error stop
+  if (ubound(pstr, 1) /= 0) error stop
+  pstr = ["01", "jk", "aq"]
+!$omp end target
+!$omp target exit data map(from: pstr)
+if (.not. associated(pstr)) error stop
+if (len(pstr) /= 2) error stop
+if (size(pstr) /= 3) error stop
+if (lbound(pstr, 1) /= -2) error stop
+if (ubound(pstr, 1) /= 0) error stop
+if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+!$omp target enter data map(tofrom: pstr2)
+!$omp target map(from: pstr2)
+  if (.not. associated(pstr2)) error stop
+  if (len(pstr2) /= 4) error stop
+  pstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: pstr2)
+if (.not. associated(pstr2)) error stop
+if (len(pstr2) /= 4) error stop
+if (pstr2 /= "HIJK") error stop
+
+
+! character(len=:), allocatable :: astr(:), astr2
+! allocate(character(len=6) :: astr(3:5))
+! allocate(character(len=8) :: astr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: astr)
+!$omp target map(from: astr)
+  if (.not. allocated(astr)) error stop
+  if (len(astr) /= 6) error stop
+  if (size(astr) /= 3) error stop
+  if (lbound(astr, 1) /= 3) error stop
+  if (ubound(astr, 1) /= 5) error stop
+  astr = ["01db45", "jk$D%S", "zutg47"]
+!$omp end target
+!$omp target exit data map(from: astr)
+if (.not. allocated(astr)) error stop
+if (len(astr) /= 6) error stop
+if (size(astr) /= 3) error stop
+if (lbound(astr, 1) /= 3) error stop
+if (ubound(astr, 1) /= 5) error stop
+if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: astr2)
+!$omp target map(from: astr2)
+  if (.not. allocated(astr2)) error stop
+  if (len(astr2) /= 8) error stop
+  astr2 = "HIJKhijk"
+!$omp end target
+!$omp target exit data map(from: astr2)
+if (.not. allocated(astr2)) error stop
+if (len(astr2) /= 8) error stop
+if (astr2 /= "HIJKhijk") error stop
+
+
+! character(len=:,kind=4), pointer :: p4str(:), p4str2
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+! allocate(character(len=5,kind=4) :: p4str2)
+
+! FAILS with value check
+
+!$omp target enter data map(tofrom: p4str)
+!$omp target map(from: p4str)
+  if (.not. associated(p4str)) error stop
+  if (len(p4str) /= 3) error stop
+  if (size(p4str) /= 3) error stop
+  if (lbound(p4str, 1) /= 2) error stop
+  if (ubound(p4str, 1) /= 4) error stop
+  p4str(:) = [4_"f85", 4_"8af", 4_"A%F"]
+!$omp end target
+!$omp target exit data map(from: p4str)
+if (.not. associated(p4str)) error stop
+if (len(p4str) /= 3) error stop
+if (size(p4str) /= 3) error stop
+if (lbound(p4str, 1) /= 2) error stop
+if (ubound(p4str, 1) /= 4) error stop
+if (p4str(2)  /= 4_"f85") error stop
+if (p4str(3)  /= 4_"8af") error stop
+if (p4str(4)  /= 4_"A%F") error stop
+
+!$omp target enter data map(tofrom: p4str2)
+!$omp target map(from: p4str2)
+  if (.not. associated(p4str2)) error stop
+  if (len(p4str2) /= 5) error stop
+  p4str2 = 4_"9875a"
+!$omp end target
+!$omp target exit data map(from: p4str2)
+if (.not. associated(p4str2)) error stop
+if (len(p4str2) /= 5) error stop
+if (p4str2 /= 4_"9875a") error stop
+
+
+! character(len=:,kind=4), allocatable :: a4str(:), a4str2
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+! allocate(character(len=9,kind=4) :: a4str2)
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+!$omp target enter data map(tofrom: a4str)
+!$omp target map(from: a4str)
+  if (.not. allocated(a4str)) error stop
+  if (len(a4str) /= 7) error stop
+  if (size(a4str) /= 6) error stop
+  if (lbound(a4str, 1) /= -2) error stop
+  if (ubound(a4str, 1) /= 3) error stop
+  ! See PR fortran/107508 why '(:)' is required
+  a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!$omp end target
+!$omp target exit data map(from: a4str)
+if (.not. allocated(a4str)) error stop
+if (len(a4str) /= 7) error stop
+if (size(a4str) /= 6) error stop
+if (lbound(a4str, 1) /= -2) error stop
+if (ubound(a4str, 1) /= 3) error stop
+if (a4str(-2) /= 4_"sf456aq") error stop
+if (a4str(-1) /= 4_"3dtzu24") error stop
+if (a4str(0)  /= 4_"_4fh7sm") error stop
+if (a4str(1)  /= 4_"=ff85s7") error stop
+if (a4str(2)  /= 4_"j=8af4d") error stop
+if (a4str(3)  /= 4_".,A%Fsz") error stop
+
+!$omp target enter data map(tofrom: a4str2)
+!$omp target map(from: a4str2)
+  if (.not. allocated(a4str2)) error stop
+  if (len(a4str2) /= 9) error stop
+  a4str2 = 4_"98765a23d"
+!$omp end target
+!$omp target exit data map(from: a4str2)
+if (.not. allocated(a4str2)) error stop
+if (len(a4str2) /= 9) error stop
+if (a4str2 /= 4_"98765a23d") error stop
+
+
+deallocate(dt%pc, dt%pc2)
+deallocate(dt%pcstr)
+deallocate(dt%pcstr2)
+
+deallocate(dt%pc4str)
+deallocate(dt%pc4str2)
+
+deallocate(ip, ip2, ia, ia2)
+deallocate(pstr)
+deallocate(pstr2)
+deallocate(astr)
+deallocate(astr2)
+
+deallocate(p4str)
+deallocate(p4str2)
+deallocate(a4str)
+deallocate(a4str2)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
new file mode 100644
index 00000000000..80d30edbfc5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
@@ -0,0 +1,392 @@
+! Check that 'map(alloc:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+  integer :: ic(2:5)
+  character(len=11) :: ccstr(3:4)
+  character(len=11,kind=4) :: cc4str(3:7)
+  integer, pointer :: pc(:)
+  character(len=:), pointer :: pcstr(:)
+  character(len=:,kind=4), pointer :: pc4str(:)
+end type t
+
+type(t) :: dt
+
+integer :: ii(5)
+character(len=11) :: clstr(-1:1)
+character(len=11,kind=4) :: cl4str(0:3)
+integer, pointer :: ip(:)
+integer, allocatable :: ia(:)
+character(len=:), pointer :: pstr(:)
+character(len=:), allocatable :: astr(:)
+character(len=:,kind=4), pointer :: p4str(:)
+character(len=:,kind=4), allocatable :: a4str(:)
+
+allocate(dt%pc(5))
+allocate(character(len=2) :: dt%pcstr(2))
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+
+allocate(ip(5), ia(8))
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=6) :: astr(3:5))
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+
+
+! integer :: ic(2:5)
+
+!$omp target enter data map(alloc: dt%ic(3:5))
+dt%ic(2) = 22
+!$omp target map(alloc: dt%ic(3:5))
+  if (size(dt%ic) /= 4) error stop
+  if (lbound(dt%ic, 1) /= 2) error stop
+  if (ubound(dt%ic, 1) /= 5) error stop
+  dt%ic(3:5) = [33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic(3:5))
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+! character(len=11) :: ccstr(3:4)
+
+!$omp target enter data map(alloc: dt%ccstr(4:4))
+dt%ccstr(3) = "12345678901"
+!$omp target map(alloc: dt%ccstr(4:4))
+  if (len(dt%ccstr) /= 11) error stop
+  if (size(dt%ccstr) /= 2) error stop
+  if (lbound(dt%ccstr, 1) /= 3) error stop
+  if (ubound(dt%ccstr, 1) /= 4) error stop
+  dt%ccstr(4:4) = ["abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr(4:4))
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7)
+
+! Value check fails
+!$omp target enter data map(alloc: dt%cc4str(4:7))
+dt%cc4str(3) = 4_"12345678901"
+!$omp target map(alloc: dt%cc4str(4:7))
+  if (len(dt%cc4str) /= 11) error stop
+  if (size(dt%cc4str) /= 5) error stop
+  if (lbound(dt%cc4str, 1) /= 3) error stop
+  if (ubound(dt%cc4str, 1) /= 7) error stop
+  dt%cc4str(4:7) = [4_"abcdefghijk", &
+               4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+               4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str(4:7))
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+! integer, pointer :: pc(:)
+! allocate(dt%pc(5))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pc(2:5))
+dt%pc(1) = 11
+!$omp target map(alloc: dt%pc(2:5))
+  if (.not. associated(dt%pc)) error stop
+  if (size(dt%pc) /= 5) error stop
+  if (lbound(dt%pc, 1) /= 1) error stop
+  if (ubound(dt%pc, 1) /= 5) error stop
+  dt%pc(2:5) = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc(2:5))
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+! character(len=:), pointer :: pcstr(:)
+! allocate(character(len=2) :: dt%pcstr(2))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+! FIXME: Disabled befause of PR108837
+!
+!!$omp target enter data map(alloc: dt%pcstr(2:2))
+!dt%pcstr(1) = "01"
+!!$omp target map(alloc: dt%pcstr(2:2))
+!  if (.not. associated(dt%pcstr)) error stop
+!  if (len(dt%pcstr) /= 2) error stop
+!  if (size(dt%pcstr) /= 2) error stop
+!  if (lbound(dt%pcstr, 1) /= 1) error stop
+!  if (ubound(dt%pcstr, 1) /= 2) error stop
+!  dt%pcstr(2:2) = ["jk"]
+!!$omp end target
+!!$omp target exit data map(from: dt%pcstr(2:2))
+!if (.not. associated(dt%pcstr)) error stop
+!if (len(dt%pcstr) /= 2) error stop
+!if (size(dt%pcstr) /= 2) error stop
+!if (lbound(dt%pcstr, 1) /= 1) error stop
+!if (ubound(dt%pcstr, 1) /= 2) error stop
+!if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:)
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+! structure element when other mapped elements from the same structure weren't mapped together with it
+
+! FIXME: Disabled befause of PR108837
+!
+!!$omp target enter data map(alloc: dt%pc4str(3:3))
+!dt%pc4str(2) = 4_"456"
+!!$omp target map(alloc: dt%pc4str(3:3))
+!  if (.not. associated(dt%pc4str)) error stop
+!  if (len(dt%pc4str) /= 3) error stop
+!  if (size(dt%pc4str) /= 2) error stop
+!  if (lbound(dt%pc4str, 1) /= 2) error stop
+!  if (ubound(dt%pc4str, 1) /= 3) error stop
+!  dt%pc4str(3:3) = [4_"tzu"]
+!!$omp end target
+!!$omp target exit data map(from: dt%pc4str(3:3))
+!if (.not. associated(dt%pc4str)) error stop
+!if (len(dt%pc4str) /= 3) error stop
+!if (size(dt%pc4str) /= 2) error stop
+!if (lbound(dt%pc4str, 1) /= 2) error stop
+!if (ubound(dt%pc4str, 1) /= 3) error stop
+!if (dt%pc4str(2) /= 4_"456") error stop
+!if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+! integer :: ii(5)
+
+!$omp target enter data map(alloc: ii(2:5))
+ii(1) = -1
+!$omp target map(alloc: ii(2:5))
+  if (size(ii) /= 5) error stop
+  if (lbound(ii, 1) /= 1) error stop
+  if (ubound(ii, 1) /= 5) error stop
+  ii(2:5) = [-2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii(2:5))
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+
+! character(len=11) :: clstr(-1:1)
+
+!$omp target enter data map(alloc: clstr(0:1))
+clstr(-1) = "12345678901"
+!$omp target map(alloc: clstr(0:1))
+  if (len(clstr) /= 11) error stop
+  if (size(clstr) /= 3) error stop
+  if (lbound(clstr, 1) /= -1) error stop
+  if (ubound(clstr, 1) /= 1) error stop
+  clstr(0:1) = ["abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr(0:1))
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+! character(len=11,kind=4) :: cl4str(0:3)
+
+!$omp target enter data map(alloc: cl4str(1:3))
+cl4str(0) = 4_"12345678901"
+!$omp target map(alloc: cl4str(1:3))
+  if (len(cl4str) /= 11) error stop
+  if (size(cl4str) /= 4) error stop
+  if (lbound(cl4str, 1) /= 0) error stop
+  if (ubound(cl4str, 1) /= 3) error stop
+  cl4str(1:3) = [4_"abcdefghijk", &
+            4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str(1:3))
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+
+! allocate(ip(5), ia(8))
+
+!$omp target enter data map(alloc: ip(2:5))
+ip(1) = 11
+!$omp target map(alloc: ip(2:5))
+  if (.not. associated(ip)) error stop
+  if (size(ip) /= 5) error stop
+  if (lbound(ip, 1) /= 1) error stop
+  if (ubound(ip, 1) /= 5) error stop
+  ip(2:5) = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip(2:5))
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+! allocate(ip(5), ia(8))
+
+!$omp target enter data map(alloc: ia(2:8))
+ia(1) = 1
+!$omp target map(alloc: ia(2:8))
+  if (.not. allocated(ia)) error stop
+  if (size(ia) /= 8) error stop
+  if (lbound(ia, 1) /= 1) error stop
+  if (ubound(ia, 1) /= 8) error stop
+  ia(2:8) = [2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia(2:8))
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+
+! character(len=:), pointer :: pstr(:)
+! allocate(character(len=2) :: pstr(-2:0))
+
+! libgomp: nvptx_alloc error: out of memory
+
+! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR
+!
+!!$omp target enter data map(alloc: pstr(-1:0))
+!pstr(-2) = "01"
+!!$omp target map(alloc: pstr(-1:0))
+!  if (.not. associated(pstr)) error stop
+!  if (len(pstr) /= 2) error stop
+!  if (size(pstr) /= 3) error stop
+!  if (lbound(pstr, 1) /= -2) error stop
+!  if (ubound(pstr, 1) /= 0) error stop
+!  pstr(-1:0) = ["jk", "aq"]
+!!$omp end target
+!!$omp target exit data map(from: pstr(-1:0))
+!if (.not. associated(pstr)) error stop
+!if (len(pstr) /= 2) error stop
+!if (size(pstr) /= 3) error stop
+!if (lbound(pstr, 1) /= -2) error stop
+!if (ubound(pstr, 1) /= 0) error stop
+!if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+
+! character(len=:), allocatable :: astr(:)
+! allocate(character(len=6) :: astr(3:5))
+
+! libgomp: nvptx_alloc error: out of memory
+
+! FIXME
+!!$omp target enter data map(alloc: astr(4:5))
+!astr(3) = "01db45"
+!!$omp target map(alloc: astr(4:5))
+!  if (.not. allocated(astr)) error stop
+!  if (len(astr) /= 6) error stop
+!  if (size(astr) /= 3) error stop
+!  if (lbound(astr, 1) /= 3) error stop
+!  if (ubound(astr, 1) /= 5) error stop
+!!!  astr(4:5) = ["jk$D%S", "zutg47"]
+!!$omp end target
+!!!$omp target exit data map(from: astr(4:5))
+!!if (.not. allocated(astr)) error stop
+!!!if (len(astr) /= 6) error stop
+!if (size(astr) /= 3) error stop
+!if (lbound(astr, 1) /= 3) error stop
+!if (ubound(astr, 1) /= 5) error stop
+!if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+!
+
+! character(len=:,kind=4), pointer :: p4str(:)
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+
+! FAILS with value check
+
+! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR
+!
+!!$omp target enter data map(alloc: p4str(3:4))
+!p4str(2) = 4_"f85"
+!!$omp target map(alloc: p4str(3:4))
+!  if (.not. associated(p4str)) error stop
+!  if (len(p4str) /= 3) error stop
+!  if (size(p4str) /= 3) error stop
+!  if (lbound(p4str, 1) /= 2) error stop
+!  if (ubound(p4str, 1) /= 4) error stop
+!  p4str(3:4) = [4_"8af", 4_"A%F"]
+!!$omp end target
+!!$omp target exit data map(from: p4str(3:4))
+!if (.not. associated(p4str)) error stop
+!if (len(p4str) /= 3) error stop
+!if (size(p4str) /= 3) error stop
+!if (lbound(p4str, 1) /= 2) error stop
+!if (ubound(p4str, 1) /= 4) error stop
+!if (p4str(2)  /= 4_"f85") error stop
+!if (p4str(3)  /= 4_"8af") error stop
+!if (p4str(4)  /= 4_"A%F") error stop
+
+! character(len=:,kind=4), allocatable :: a4str(:)
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+! FIXME: Disabled befause of PR108838
+!!$omp target enter data map(alloc: a4str(-1:3))
+!!a4str(-2) = 4_"sf456aq"
+!!$omp target map(alloc: a4str(-1:3))
+!  if (.not. allocated(a4str)) error stop
+!  if (len(a4str) /= 7) error stop
+!  if (size(a4str) /= 6) error stop
+!  if (lbound(a4str, 1) /= -2) error stop
+!  if (ubound(a4str, 1) /= 3) error stop
+!  a4str(-1:3) = [4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!!$omp end target
+!!$omp target exit data map(from: a4str(-1:3))
+!if (.not. allocated(a4str)) error stop
+!if (len(a4str) /= 7) error stop
+!if (size(a4str) /= 6) error stop
+!if (lbound(a4str, 1) /= -2) error stop
+!if (ubound(a4str, 1) /= 3) error stop
+!if (a4str(-2) /= 4_"sf456aq") error stop
+!if (a4str(-1) /= 4_"3dtzu24") error stop
+!if (a4str(0)  /= 4_"_4fh7sm") error stop
+!if (a4str(1)  /= 4_"=ff85s7") error stop
+!if (a4str(2)  /= 4_"j=8af4d") error stop
+!if (a4str(3)  /= 4_".,A%Fsz") error stop
+
+deallocate(dt%pc)
+deallocate(dt%pcstr)
+
+deallocate(dt%pc4str)
+
+deallocate(ip, ia)
+deallocate(pstr)
+deallocate(astr)
+
+deallocate(p4str)
+deallocate(a4str)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90
new file mode 100644
index 00000000000..f129f559336
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90
@@ -0,0 +1,78 @@
+module m
+  implicit none
+  character(len=:), allocatable :: strA(:), strA2
+  character(len=:), pointer :: strP(:), strP2
+  !$omp declare target enter(strA,strA2,strP,strP2)
+contains
+  subroutine opt_map(str1, str2, str3)
+    character(len=:), allocatable :: str1, str2, str3, str4
+    optional :: str2, str3 
+
+    if (.not.present(str2)) error stop
+    if (present(str3))  error stop
+
+    !$omp target map(str1,str2,str3,str4)
+      if (allocated(str1)) error stop
+      if (allocated(str2)) error stop
+      if (present(str3)) error stop
+      if (allocated(str4)) error stop
+    !$omp end target
+  end
+  subroutine call_opt()
+    character(len=:), allocatable :: str1, str2
+    call opt_map(str1, str2)
+  end
+  subroutine test
+   !$omp declare target
+   if (.not. allocated(strA)) error stop
+   !if (.not. allocated(strA2)) error stop
+   if (.not. associated(strP)) error stop
+   !if (.not. associated(strP2)) error stop
+
+    ! ensure length was updated as well
+    if (len(strA) /= 3) error stop
+    if (len(strA2) /= 5) error stop
+    if (len(strP) /= 4) error stop
+    if (len(strP2) /= 8) error stop
+!    if (any (strA /= ['Hav', 'e f', 'un!'])) error stop
+!    if (strA2 /= 'Hello') error stop
+!    if (any (strP /= ['abcd', 'efgh', 'ijkl'])) error stop
+!    if (strP2 /= 'TestCase') error stop
+!
+!    strA = ['123', '456', '789']
+!    strA2 = 'World'
+!    strP = ['ABCD', 'EFGH', 'IJKL']
+!    strP2 = 'Passed!!'
+  end
+end
+
+program main
+  use m
+  implicit none
+  call call_opt
+
+  strA = ['Hav', 'e f', 'un!']
+  strA2 = 'Hello'
+  allocate(character(len=4) :: strP(3))
+  strP = ['abcd', 'efgh', 'ijkl']
+  allocate(character(len=8) :: strP2)
+  strP2 = 'TestCase'
+
+  !$omp target enter data map(always, to: strA, strA2)
+  !$omp target enter data map(to: strP, strP2)
+  !$omp target
+    call test()
+  !$omp end target
+  !$omp target exit data map(always, from: strA, strA2, strP, strP2)
+
+  if (len(strA) /= 3) error stop
+  if (len(strA2) /= 5) error stop
+  if (len(strP) /= 4) error stop
+  if (len(strP2) /= 8) error stop
+!  if (any (strA /= ['123', '456', '789'])) error stop
+!  if (strA2 /= 'World') error stop
+!  if (any(strP /= ['ABCD', 'EFGH', 'IJKL'])) error stop
+!  if (strP2 /= 'Passed!!') error stop
+
+!  deallocate(strP, strP2, strA, strA2)
+end