[OpenMP, Fortran] Add structure/derived-type element mapping

gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clauses): Match also derived-type
	component refs in OMP_CLAUSE_MAP.
	(resolve_omp_clauses): Resolve those.
	* trans-openmp.c (gfc_trans_omp_array_section, gfc_trans_omp_clauses):
	Handle OpenMP structure-element mapping.
	(gfc_trans_oacc_construct, gfc_trans_oacc_executable_directive,
	(gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Update
	add openacc=true in gfc_trans_omp_clauses call.

gcc/testsuite/ChangeLog:

	* gfortran.dg/goacc/finalize-1.f: Update dump scan pattern.
	* gfortran.dg/gomp/map-1.f90: Update dg-error.
	* gfortran.dg/gomp/map-2.f90: New test.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/struct-elem-map-1.f90: New test.

6 files changed, 595 insertions, 118 deletions

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index e681903c7c2..7de2f6e1b1d 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1464,7 +1464,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      head = NULL;
 	      if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
 					       false, NULL, &head,
-					       true) == MATCH_YES)
+					       true, true) == MATCH_YES)
 		{
 		  gfc_omp_namelist *n;
 		  for (n = *head; n; n = n->next)
@@ -4553,7 +4553,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 
 		    /* Look through component refs to find last array
 		       reference.  */
-		    if (openacc && resolved)
+		    if (resolved)
 		      {
 			/* The "!$acc cache" directive allows rectangular
 			   subarrays to be specified, with some restrictions
@@ -4563,6 +4563,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			   arr(-n:n,-n:n) could be contiguous even if it looks
 			   like it may not be.  */
 			if (list != OMP_LIST_CACHE
+			    && list != OMP_LIST_DEPEND
 			    && !gfc_is_simply_contiguous (n->expr, false, true)
 			    && gfc_is_not_contiguous (n->expr))
 			  gfc_error ("Array is not contiguous at %L",
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index b2645e723d5..f35ff4191c2 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2092,10 +2092,11 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
 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)
+			     tree &node, tree &node2, tree &node3, tree &node4)
 {
   gfc_se se;
   tree ptr, ptr2;
+  tree elemsz = NULL_TREE;
 
   gfc_init_se (&se, NULL);
 
@@ -2104,7 +2105,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
       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 (ptr));
+      OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
+      elemsz = OMP_CLAUSE_SIZE (node);
     }
   else
     {
@@ -2114,14 +2116,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
       gfc_add_block_to_block (block, &se.pre);
       OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
 						    GFC_TYPE_ARRAY_RANK (type));
-      tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+      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);
     }
-  gfc_add_block_to_block (block, &se.post);
+  gcc_assert (se.post.head == NULL_TREE);
   ptr = fold_convert (build_pointer_type (char_type_node), ptr);
   OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+  ptr = fold_convert (ptrdiff_type_node, ptr);
 
   if (POINTER_TYPE_P (TREE_TYPE (decl))
       && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
@@ -2134,28 +2137,71 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
       OMP_CLAUSE_SIZE (node4) = size_int (0);
       decl = build_fold_indirect_ref (decl);
     }
-  ptr = fold_convert (sizetype, ptr);
+  else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER
+	   && n->expr->ts.type == BT_CHARACTER
+	   && 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)
+	map_kind = OMP_CLAUSE_MAP_KIND (node);
+      else
+	map_kind = GOMP_MAP_ALLOC;
+      gcc_assert (se.string_length);
+      node4 = build_omp_clause (input_location, 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);
+    }
   if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
     {
+      tree desc_node;
       tree type = TREE_TYPE (decl);
       ptr2 = gfc_conv_descriptor_data_get (decl);
-      node2 = build_omp_clause (input_location,
-				OMP_CLAUSE_MAP);
-      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
-      OMP_CLAUSE_DECL (node2) = decl;
-      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+      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)
+	{
+	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
+	  node2 = node;
+	  node = desc_node;  /* Needs to come first.  */
+	}
+      else
+	{
+	  OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
+	  node2 = desc_node;
+	}
       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);
+      /* This purposely does not include GOMP_MAP_ALWAYS_POINTER.  The extra
+	 cast prevents gimplify.c from recognising it as being part of the
+	 struct – and adding an 'alloc: for the 'desc.data' pointer, which
+	 would break as the 'desc' (the descriptor) is also mapped
+	 (see node4 above).  */
       if (ptr_kind == GOMP_MAP_ATTACH_DETACH)
 	STRIP_NOPS (OMP_CLAUSE_DECL (node3));
     }
   else
     {
       if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
-	ptr2 = build_fold_addr_expr (decl);
+	{
+	  tree offset;
+	  ptr2 = build_fold_addr_expr (decl);
+	  offset = fold_build2 (MINUS_EXPR, ptrdiff_type_node, ptr,
+				fold_convert (ptrdiff_type_node, ptr2));
+	  offset = build2 (TRUNC_DIV_EXPR, ptrdiff_type_node,
+			   offset, fold_convert (ptrdiff_type_node, elemsz));
+	  offset = build4_loc (input_location, ARRAY_REF,
+			       TREE_TYPE (TREE_TYPE (decl)),
+			       decl, offset, NULL_TREE, NULL_TREE);
+	  OMP_CLAUSE_DECL (node) = offset;
+	}
       else
 	{
 	  gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
@@ -2166,14 +2212,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
       OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
       OMP_CLAUSE_DECL (node3) = decl;
     }
-  ptr2 = fold_convert (sizetype, ptr2);
-  OMP_CLAUSE_SIZE (node3)
-    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+  ptr2 = fold_convert (ptrdiff_type_node, ptr2);
+  OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
+					 ptr, ptr2);
 }
 
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
-		       locus where, bool declare_simd = false)
+		       locus where, bool declare_simd = false,
+		       bool openacc = false)
 {
   tree omp_clauses = NULL_TREE, chunk_size, c;
   int list, ifc;
@@ -2488,6 +2535,67 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      tree node2 = NULL_TREE;
 	      tree node3 = NULL_TREE;
 	      tree node4 = NULL_TREE;
+
+	      switch (n->u.map_op)
+		{
+		case OMP_MAP_ALLOC:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+		  break;
+		case OMP_MAP_IF_PRESENT:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+		  break;
+		case OMP_MAP_ATTACH:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
+		  break;
+		case OMP_MAP_TO:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
+		  break;
+		case OMP_MAP_FROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
+		  break;
+		case OMP_MAP_TOFROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
+		  break;
+		case OMP_MAP_ALWAYS_TO:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
+		  break;
+		case OMP_MAP_ALWAYS_FROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
+		  break;
+		case OMP_MAP_ALWAYS_TOFROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
+		  break;
+		case OMP_MAP_RELEASE:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
+		  break;
+		case OMP_MAP_DELETE:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
+		  break;
+		case OMP_MAP_DETACH:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
+		  break;
+		case OMP_MAP_FORCE_ALLOC:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
+		  break;
+		case OMP_MAP_FORCE_TO:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
+		  break;
+		case OMP_MAP_FORCE_FROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
+		  break;
+		case OMP_MAP_FORCE_TOFROM:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
+		  break;
+		case OMP_MAP_FORCE_PRESENT:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
+		  break;
+		case OMP_MAP_FORCE_DEVICEPTR:
+		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
+		  break;
+		default:
+		  gcc_unreachable ();
+		}
+
 	      tree decl = gfc_trans_omp_variable (n->sym, false);
 	      if (DECL_P (decl))
 		TREE_ADDRESSABLE (decl) = 1;
@@ -2496,7 +2604,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      && n->expr->ref->u.ar.type == AR_FULL))
 		{
 		  tree present = gfc_omp_check_optional_argument (decl, true);
-		  if (n->sym->ts.type == BT_CLASS)
+		  if (openacc && n->sym->ts.type == BT_CLASS)
 		    {
 		      tree type = TREE_TYPE (decl);
 		      if (n->sym->attr.optional)
@@ -2724,8 +2832,42 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      /* Last component is a scalar.  */
 		      gfc_conv_expr (&se, n->expr);
 		      gfc_add_block_to_block (block, &se.pre);
-		      OMP_CLAUSE_DECL (node) = se.expr;
+		      /* For BT_CHARACTER a pointer is returned.  */
+		      OMP_CLAUSE_DECL (node)
+			= POINTER_TYPE_P (TREE_TYPE (se.expr))
+			  ? build_fold_indirect_ref (se.expr) : se.expr;
 		      gfc_add_block_to_block (block, &se.post);
+		      if (sym_attr.pointer || sym_attr.allocatable)
+			{
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node2,
+						   openacc
+						   ? GOMP_MAP_ATTACH_DETACH
+						   : GOMP_MAP_ALWAYS_POINTER);
+			  OMP_CLAUSE_DECL (node2)
+			    = POINTER_TYPE_P (TREE_TYPE (se.expr))
+			      ? se.expr :  gfc_build_addr_expr (NULL, se.expr);
+			  OMP_CLAUSE_SIZE (node2) = size_int (0);
+			  if (!openacc
+			      && n->expr->ts.type == BT_CHARACTER
+			      && n->expr->ts.deferred)
+			    {
+			      gcc_assert (se.string_length);
+			      tree tmp = gfc_get_char_type (n->expr->ts.kind);
+			      OMP_CLAUSE_SIZE (node)
+				= fold_build2 (MULT_EXPR, size_type_node,
+					       fold_convert (size_type_node,
+							     se.string_length),
+					       TYPE_SIZE_UNIT (tmp));
+			      node3 = build_omp_clause (input_location,
+							OMP_CLAUSE_MAP);
+			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO);
+			      OMP_CLAUSE_DECL (node3) = se.string_length;
+			      OMP_CLAUSE_SIZE (node3)
+				= TYPE_SIZE_UNIT (gfc_charlen_type_node);
+			    }
+			}
 		      goto finalize_map_clause;
 		    }
 
@@ -2752,7 +2894,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  if (lastcomp->u.c.component->ts.type == BT_DERIVED
 		      || lastcomp->u.c.component->ts.type == BT_CLASS)
 		    {
-		      if (sym_attr.allocatable || sym_attr.pointer)
+		      if (sym_attr.pointer || (openacc && sym_attr.allocatable))
 			{
 			  tree data, size;
 
@@ -2773,7 +2915,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  node2 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
 			  OMP_CLAUSE_SET_MAP_KIND (node2,
-						   GOMP_MAP_ATTACH_DETACH);
+						   openacc
+						   ? GOMP_MAP_ATTACH_DETACH
+						   : GOMP_MAP_ALWAYS_POINTER);
 			  OMP_CLAUSE_DECL (node2) = data;
 			  OMP_CLAUSE_SIZE (node2) = size_int (0);
 			}
@@ -2800,32 +2944,82 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 
 		      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
 			{
+			  gomp_map_kind map_kind;
+			  tree desc_node;
 			  tree type = TREE_TYPE (inner);
 			  tree ptr = gfc_conv_descriptor_data_get (inner);
 			  ptr = build_fold_indirect_ref (ptr);
 			  OMP_CLAUSE_DECL (node) = ptr;
-			  node2 = build_omp_clause (input_location,
-						    OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
-			  OMP_CLAUSE_DECL (node2) = inner;
-			  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-			  node3 = build_omp_clause (input_location,
-						    OMP_CLAUSE_MAP);
-			  OMP_CLAUSE_SET_MAP_KIND (node3,
-						   GOMP_MAP_ATTACH_DETACH);
-			  OMP_CLAUSE_DECL (node3)
-			    = gfc_conv_descriptor_data_get (inner);
-			  STRIP_NOPS (OMP_CLAUSE_DECL (node3));
-			  OMP_CLAUSE_SIZE (node3) = size_int (0);
 			  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));
+			  if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
+			    map_kind = 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
+			    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 (input_location,
+							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);
+			  desc_node = build_omp_clause (input_location,
+							OMP_CLAUSE_MAP);
+			  if (openacc)
+			    OMP_CLAUSE_SET_MAP_KIND (desc_node,
+						     GOMP_MAP_TO_PSET);
+			  else
+			    OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
+			  OMP_CLAUSE_DECL (desc_node) = inner;
+			  OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
+			  if (openacc)
+			    node2 = desc_node;
+			  else
+			    {
+			      node2 = node;
+			      node = desc_node;  /* Put first.  */
+			    }
+			  node3 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node3,
+						   openacc
+						   ? GOMP_MAP_ATTACH_DETACH
+						   : GOMP_MAP_ALWAYS_POINTER);
+			  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);
 			}
 		      else
 			OMP_CLAUSE_DECL (node) = inner;
@@ -2837,9 +3031,11 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			   && lastcomp->next->type == REF_ARRAY
 			   && lastcomp->next->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,
-						   GOMP_MAP_ATTACH_DETACH,
-						   node, node2, node3, node4);
+						   kind, node, node2, node3,
+						   node4);
 		    }
 		}
 	      else  /* An array element or array section.  */
@@ -2851,65 +3047,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		}
 
 	      finalize_map_clause:
-	      switch (n->u.map_op)
-		{
-		case OMP_MAP_ALLOC:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
-		  break;
-		case OMP_MAP_IF_PRESENT:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
-		  break;
-		case OMP_MAP_ATTACH:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
-		  break;
-		case OMP_MAP_TO:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
-		  break;
-		case OMP_MAP_FROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
-		  break;
-		case OMP_MAP_TOFROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
-		  break;
-		case OMP_MAP_ALWAYS_TO:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
-		  break;
-		case OMP_MAP_ALWAYS_FROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
-		  break;
-		case OMP_MAP_ALWAYS_TOFROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
-		  break;
-		case OMP_MAP_RELEASE:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
-		  break;
-		case OMP_MAP_DELETE:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
-		  break;
-		case OMP_MAP_DETACH:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
-		  break;
-		case OMP_MAP_FORCE_ALLOC:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
-		  break;
-		case OMP_MAP_FORCE_TO:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
-		  break;
-		case OMP_MAP_FORCE_FROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
-		  break;
-		case OMP_MAP_FORCE_TOFROM:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
-		  break;
-		case OMP_MAP_FORCE_PRESENT:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
-		  break;
-		case OMP_MAP_FORCE_DEVICEPTR:
-		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
-		  break;
-		default:
-		  gcc_unreachable ();
-		}
+
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	      if (node2)
 		omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
@@ -3661,7 +3799,7 @@ gfc_trans_oacc_construct (gfc_code *code)
 
   gfc_start_block (&block);
   oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-					code->loc);
+					code->loc, false, true);
   stmt = gfc_trans_omp_code (code->block->next, true);
   stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
 		     oacc_clauses);
@@ -3697,7 +3835,7 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
 
   gfc_start_block (&block);
   oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-					code->loc);
+					code->loc, false, true);
   stmt = build1_loc (input_location, construct_code, void_type_node, 
 		     oacc_clauses);
   gfc_add_expr_to_block (&block, stmt);
@@ -4522,7 +4660,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
       if (construct_code == OACC_KERNELS)
 	construct_clauses.lists[OMP_LIST_REDUCTION] = NULL;
       oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
-					    code->loc);
+					    code->loc, false, true);
     }
   if (!loop_clauses.seq)
     pblock = &block;
@@ -5703,7 +5841,7 @@ gfc_trans_oacc_declare (gfc_code *code)
   gfc_start_block (&block);
 
   oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
-					code->loc);
+					code->loc, false, true);
   stmt = gfc_trans_omp_code (code->block->next, true);
   stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
 		     oacc_clauses);
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index ca642156e9f..74fa408082b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -20,7 +20,7 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_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:\\*\\(c_char \\*\\) 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: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) 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: \\(integer\\(kind=8\\)\\) parm\\.0\\.data - \\(integer\\(kind=8\\)\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
@@ -32,6 +32,6 @@
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm\\.1\\.data - \\(integer\\(kind=8\\)\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-1.f90 b/gcc/testsuite/gfortran.dg/gomp/map-1.f90
index e78b56c8f39..831feffcc43 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-1.f90
@@ -57,18 +57,20 @@ subroutine test(aas)
   !$omp target map(j(:))
   !$omp end target
 
-  !$omp target map(j(1:9:2)) ! { dg-error "Stride should not be specified for array section in MAP clause" }
+  !$omp target map(j(1:9:2))
+ ! { dg-error "Array is not contiguous" "" { target *-*-* } 60 }
+ ! { dg-error "Stride should not be specified for array section in MAP clause" "" { target *-*-* } 60 }
   !$omp end target
 
   !$omp target map(aas(5:))
   !$omp end target
-  ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 63 }
-  ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 63 }
+  ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 65 }
+  ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 65 }
 
   !$omp target map(aas(:))
   !$omp end target
-  ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 68 }
-  ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 68 }
+  ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 70 }
+  ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 70 }
 
   !$omp target map(aas) ! { dg-error "Assumed size array" }
   !$omp end target
@@ -81,29 +83,28 @@ subroutine test(aas)
 
   !$omp target map(k(5:))
   !$omp end target
-  ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 82 }
-  ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 82 }
+  ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 84 }
+  ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 84 }
 
   !$omp target map(k(5:,:,3))
   !$omp end target
-  ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 87 }
-  ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 87 }
+  ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 89 }
+  ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 89 }
 
   !$omp target map(tt)
   !$omp end target
 
-  !$omp target map(tt%i) ! { dg-error "Syntax error in OpenMP variable list" }
+  !$omp target map(tt%k) ! { dg-error "not a member of" }
   !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
 
-  !$omp target map(tt%j) ! { dg-error "Syntax error in OpenMP variable list" }
-  !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+  !$omp target map(tt%j)
+  !$omp end target
 
-  ! broken test
-  !$omp target map(tt%j(1)) ! { dg-error "Syntax error in OpenMP variable list" }
-  !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+  !$omp target map(tt%j(1))
+  !$omp end target
 
-  !$omp target map(tt%j(1:)) ! { dg-error "Syntax error in OpenMP variable list" }
-  !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+  !$omp target map(tt%j(1:))
+  !$omp end target
 
   !$omp target map(tp) ! { dg-error "THREADPRIVATE object 'tp' in MAP clause" }
   !$omp end target
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
new file mode 100644
index 00000000000..73c4f5a87d0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
@@ -0,0 +1,6 @@
+type t
+  integer :: i
+end type t
+type(t) v
+!$omp target enter data map(to:v%i, v%i)  ! { dg-error "appears more than once in map clauses" }
+end
diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
new file mode 100644
index 00000000000..f18eeb90165
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
@@ -0,0 +1,331 @@
+! { dg-do run }
+!
+! Test OpenMP 4.5 structure-element mapping
+
+! TODO: character(kind=4,...) needs to be tested, but depends on
+!       PR fortran/95837
+! TODO: ...%str4 should be tested but that currently fails due to
+!       PR fortran/95868 (see commented lined)
+! TODO: Test also array-valued var, nested derived types,
+!       type-extended types.
+
+program main
+  implicit none
+
+  type t2
+    integer :: a, b
+    ! For complex, assume small integers are exactly representable
+    complex(kind=8) :: c
+    integer :: d(10)
+    integer, pointer :: e => null(), f(:) => null()
+    character(len=5) :: str1
+    character(len=5) :: str2(4)
+    character(len=:), pointer :: str3 => null()
+    character(len=:), pointer :: str4(:) => null()
+  end type t2
+
+  integer :: i
+
+  call one ()
+  call two ()
+  call three ()
+  call four ()
+  call five ()
+  call six ()
+  call seven ()
+  call eight ()
+
+contains
+  ! Implicitly mapped – but no pointers are mapped
+  subroutine one() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "one" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%e, source=99)
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str3, source="HelloWorld")
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom:var)
+      if (var%a /= 1) stop 1
+      if (var%b /= 2)  stop 2
+      if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+      if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+      if (var%str1 /= "abcde") stop 5
+      if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+    !$omp end target
+
+    deallocate(var%e, var%f, var%str3, var%str4)
+  end subroutine one
+
+  ! Explicitly mapped – all and full arrays
+  subroutine two() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "two" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%e, source=99)
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str3, source="HelloWorld")
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom: var%a, var%b, var%c, var%d, var%e, var%f, &
+    !$omp&                   var%str1, var%str2, var%str3, var%str4)
+      if (var%a /= 1) stop 1
+      if (var%b /= 2)  stop 2
+      if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+      if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+      if (var%str1 /= "abcde") stop 5
+      if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+
+      if (.not. associated (var%e)) stop 7
+      if (var%e /= 99) stop 8
+      if (.not. associated (var%f)) stop 9
+      if (size (var%f) /= 4) stop 10
+      if (any (var%f /= [22, 33, 44, 55])) stop 11
+      if (.not. associated (var%str3)) stop 12
+      if (len (var%str3) /= len ("HelloWorld")) stop 13
+      if (var%str3 /= "HelloWorld") stop 14
+      if (.not. associated (var%str4)) stop 15
+      if (len (var%str4) /= 5) stop 16
+      if (size (var%str4) /= 2) stop 17
+      if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
+    !$omp end target
+
+    deallocate(var%e, var%f, var%str3, var%str4)
+  end subroutine two
+
+  ! Explicitly mapped – one by one but full arrays
+  subroutine three() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "three" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%e, source=99)
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str3, source="HelloWorld")
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom: var%a)
+      if (var%a /= 1) stop 1
+    !$omp end target
+    !$omp target map(tofrom: var%b)
+      if (var%b /= 2)  stop 2
+    !$omp end target
+    !$omp target map(tofrom: var%c)
+      if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+    !$omp end target
+    !$omp target map(tofrom: var%d)
+      if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+    !$omp end target
+    !$omp target map(tofrom: var%str1)
+      if (var%str1 /= "abcde") stop 5
+    !$omp end target
+    !$omp target map(tofrom: var%str2)
+      if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+    !$omp end target
+
+    !$omp target map(tofrom: var%e)
+      if (.not. associated (var%e)) stop 7
+      if (var%e /= 99) stop 8
+    !$omp end target
+    !$omp target map(tofrom: var%f)
+      if (.not. associated (var%f)) stop 9
+      if (size (var%f) /= 4) stop 10
+      if (any (var%f /= [22, 33, 44, 55])) stop 11
+    !$omp end target
+    !$omp target map(tofrom: var%str3)
+      if (.not. associated (var%str3)) stop 12
+      if (len (var%str3) /= len ("HelloWorld")) stop 13
+      if (var%str3 /= "HelloWorld") stop 14
+    !$omp end target
+    !$omp target map(tofrom: var%str4)
+      if (.not. associated (var%str4)) stop 15
+      if (len (var%str4) /= 5) stop 16
+      if (size (var%str4) /= 2) stop 17
+      if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
+    !$omp end target
+
+    deallocate(var%e, var%f, var%str3, var%str4)
+  end subroutine three
+
+  ! Explicitly mapped – all but only subarrays
+  subroutine four() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "four" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+!   !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3), var%str4(2:2))
+    !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3))
+      if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+
+      if (.not. associated (var%f)) stop 9
+      if (size (var%f) /= 4) stop 10
+      if (any (var%f(2:3) /= [33, 44])) stop 11
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+    !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine four
+
+  ! Explicitly mapped – all but only subarrays and one by one
+  subroutine five() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "five" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom: var%d(4:7))
+      if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+    !$omp end target
+    !$omp target map(tofrom: var%str2(2:3))
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+    !$omp end target
+
+    !$omp target map(tofrom: var%f(2:3))
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (any (var%f(2:3) /= [33, 44])) stop 11
+    !$omp end target
+!  !$omp target map(tofrom: var%str4(2:2))
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+!  !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine five
+
+  ! Explicitly mapped – all but only array elements
+  subroutine six() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "six" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+!   !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), var%str4(2))
+    !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3))
+      if (var%d(5) /= -3*5) stop 4
+      if (var%str2(3) /= "ABCDE") stop 6
+
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (var%f(3) /= 44) stop 11
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+    !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine six
+
+  ! Explicitly mapped – all but only array elements and one by one
+  subroutine seven() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "seven" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+    allocate (var%f, source=[22, 33, 44, 55])
+    allocate (var%str4, source=["Let's", "Go!!!"])
+
+    !$omp target map(tofrom: var%d(5))
+      if (var%d(5) /= (-3*5)) stop 4
+    !$omp end target
+    !$omp target map(tofrom: var%str2(2:3))
+      if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+    !$omp end target
+
+    !$omp target map(tofrom: var%f(2:3))
+     if (.not. associated (var%f)) stop 9
+     if (size (var%f) /= 4) stop 10
+     if (any (var%f(2:3) /= [33, 44])) stop 11
+    !$omp end target
+!   !$omp target map(tofrom: var%str4(2:2))
+!     if (.not. associated (var%str4)) stop 15
+!     if (len (var%str4) /= 5) stop 16
+!     if (size (var%str4) /= 2) stop 17
+!     if (var%str4(2) /= "Go!!!") stop 18
+!   !$omp end target
+
+    deallocate(var%f, var%str4)
+  end subroutine seven
+
+  ! Check mapping of NULL pointers
+  subroutine eight() 
+    type(t2) :: var, var2(4)
+    type(t2), pointer :: var3, var4(:)
+
+    print '(g0)', '==== TESTCASE "eight" ===='
+
+    var = t2(a = 1, &
+             b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+             d = [(-3*i, i = 1, 10)], &
+             str1 = "abcde", &
+             str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+
+!    !$omp target map(tofrom: var%e, var%f, var%str3, var%str4)
+    !$omp target map(tofrom: var%e, var%str3)
+      if (associated (var%e)) stop 1
+!      if (associated (var%f)) stop 2
+      if (associated (var%str3)) stop 3
+!      if (associated (var%str4)) stop 4
+    !$omp end target
+  end subroutine eight
+
+end program main