OpenMP: Expand "declare mapper" mappers for target {enter,exit,} data directives

This patch allows 'declare mapper' mappers to be used on 'omp target
data', 'omp target enter data' and 'omp target exit data' directives.
For each of these, only explicit mappings are supported, unlike for
'omp target' directives where implicit uses of variables inside an
offload region might trigger mappers also.

Each of C, C++ and Fortran are supported.

The patch also adjusts 'map kind decay' to match OpenMP 5.2 semantics,
which is particularly important with regard to 'exit data' operations.

2023-07-06  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
	C_ORT_OMP_EXIT_DATA.
	(c_omp_instantiate_mappers): Add region type parameter.
	* c-omp.cc (omp_split_map_kind, omp_join_map_kind,
	omp_map_decayed_kind): New functions.
	(omp_instantiate_mapper): Add ORT parameter.  Implement map kind decay
	for instantiated mapper clauses.
	(c_omp_instantiate_mappers): Add ORT parameter, pass to
	omp_instantiate_mapper.

gcc/c/
	* c-parser.cc (c_parser_omp_target_data): Instantiate mappers for
	'omp target data'.
	(c_parser_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(c_parser_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.
	(c_parser_omp_target): Add c_omp_region_type argument to
	c_omp_instantiate_mappers call.
	* c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype.

gcc/cp/
	* parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp
	target data'.
	(cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.
	(cp_parser_omp_target): Add c_omp_region_type argument to
	c_omp_instantiate_mappers call.
	* pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other
	than just C_ORT_OMP_TARGET.
	(tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE,
	OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza.
	* semantics.cc (cxx_omp_map_array_section): Avoid calling
	build_array_ref for non-array/non-pointer bases (error reported
	already).

gcc/fortran/
	* trans-openmp.cc (omp_split_map_op, omp_join_map_op,
	omp_map_decayed_kind): New functions.
	(gfc_trans_omp_instantiate_mapper): Add CD parameter.  Implement map
	kind decay.
	(gfc_trans_omp_instantiate_mappers): Add CD parameter.  Pass to above
	function.
	(gfc_trans_omp_target_data): Instantiate mappers for 'omp target data'.
	(gfc_trans_omp_target_enter_data): Instantiate mappers for 'omp target
	enter data'.
	(gfc_trans_omp_target_exit_data): Instantiate mappers for 'omp target
	exit data'.

gcc/testsuite/
	* c-c++-common/gomp/declare-mapper-15.c: New test.
	* c-c++-common/gomp/declare-mapper-16.c: New test.
	* g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output.
	* gfortran.dg/gomp/declare-mapper-22.f90: New test.
	* gfortran.dg/gomp/declare-mapper-23.f90: New test.

18 files changed, 657 insertions, 36 deletions

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 0038eba..914ac60 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,15 @@
+2023-07-06  Julian Brown  <julian@codesourcery.com>
+
+	* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
+	C_ORT_OMP_EXIT_DATA.
+	(c_omp_instantiate_mappers): Add region type parameter.
+	* c-omp.cc (omp_split_map_kind, omp_join_map_kind,
+	omp_map_decayed_kind): New functions.
+	(omp_instantiate_mapper): Add ORT parameter.  Implement map kind decay
+	for instantiated mapper clauses.
+	(c_omp_instantiate_mappers): Add ORT parameter, pass to
+	omp_instantiate_mapper.
+
 2023-06-30  Julian Brown  <julian@codesourcery.com>
 
 	* c-common.h (omp_mapper_list): Add forward declaration.
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index f970e6d..6c905ba 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1322,7 +1322,7 @@ extern void c_omp_mark_declare_variant (location_t, tree, tree);
 extern void c_omp_adjust_map_clauses (tree, bool);
 template<typename T> struct omp_mapper_list;
 extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree);
-extern tree c_omp_instantiate_mappers (tree);
+extern tree c_omp_instantiate_mappers (tree, enum c_omp_region_type);
 
 namespace omp_addr_tokenizer { struct omp_addr_token; }
 typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index b4cae46..390d6ae 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4360,13 +4360,189 @@ remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
   return NULL_TREE;
 }
 
+static enum gomp_map_kind
+omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p,
+		    bool *present_p)
+{
+  *force_p = *always_p = *present_p = false;
+
+  switch (op)
+    {
+    case GOMP_MAP_FORCE_ALLOC:
+    case GOMP_MAP_FORCE_TO:
+    case GOMP_MAP_FORCE_FROM:
+    case GOMP_MAP_FORCE_TOFROM:
+    case GOMP_MAP_FORCE_PRESENT:
+      *force_p = true;
+      break;
+    case GOMP_MAP_ALWAYS_TO:
+    case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_ALWAYS_TOFROM:
+      *always_p = true;
+      break;
+    case GOMP_MAP_ALWAYS_PRESENT_TO:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+      *always_p = true;
+      /* Fallthrough.  */
+    case GOMP_MAP_PRESENT_ALLOC:
+    case GOMP_MAP_PRESENT_TO:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_PRESENT_TOFROM:
+      *present_p = true;
+      break;
+    default:
+      ;
+    }
+
+  switch (op)
+    {
+    case GOMP_MAP_ALLOC:
+    case GOMP_MAP_FORCE_ALLOC:
+    case GOMP_MAP_PRESENT_ALLOC:
+      return GOMP_MAP_ALLOC;
+    case GOMP_MAP_TO:
+    case GOMP_MAP_FORCE_TO:
+    case GOMP_MAP_ALWAYS_TO:
+    case GOMP_MAP_PRESENT_TO:
+    case GOMP_MAP_ALWAYS_PRESENT_TO:
+      return GOMP_MAP_TO;
+    case GOMP_MAP_FROM:
+    case GOMP_MAP_FORCE_FROM:
+    case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+      return GOMP_MAP_FROM;
+    case GOMP_MAP_TOFROM:
+    case GOMP_MAP_FORCE_TOFROM:
+    case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_PRESENT_TOFROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+      return GOMP_MAP_TOFROM;
+    default:
+      ;
+    }
+
+  return op;
+}
+
+static enum gomp_map_kind
+omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p,
+		   bool present_p)
+{
+  gcc_assert (!force_p || !(always_p || present_p));
+
+  switch (op)
+    {
+    case GOMP_MAP_ALLOC:
+      if (force_p)
+	return GOMP_MAP_FORCE_ALLOC;
+      else if (present_p)
+	return GOMP_MAP_PRESENT_ALLOC;
+      break;
+
+    case GOMP_MAP_TO:
+      if (force_p)
+	return GOMP_MAP_FORCE_TO;
+      else if (always_p && present_p)
+	return GOMP_MAP_ALWAYS_PRESENT_TO;
+      else if (always_p)
+	return GOMP_MAP_ALWAYS_TO;
+      else if (present_p)
+	return GOMP_MAP_PRESENT_TO;
+      break;
+
+    case GOMP_MAP_FROM:
+      if (force_p)
+	return GOMP_MAP_FORCE_FROM;
+      else if (always_p && present_p)
+	return GOMP_MAP_ALWAYS_PRESENT_FROM;
+      else if (always_p)
+	return GOMP_MAP_ALWAYS_FROM;
+      else if (present_p)
+	return GOMP_MAP_PRESENT_FROM;
+      break;
+
+    case GOMP_MAP_TOFROM:
+      if (force_p)
+	return GOMP_MAP_FORCE_TOFROM;
+      else if (always_p && present_p)
+	return GOMP_MAP_ALWAYS_PRESENT_TOFROM;
+      else if (always_p)
+	return GOMP_MAP_ALWAYS_TOFROM;
+      else if (present_p)
+	return GOMP_MAP_PRESENT_TOFROM;
+      break;
+
+    default:
+      ;
+    }
+
+  return op;
+}
+
+/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive").  Return the
+   map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS
+   specified on the clause that invokes the mapper.  See also
+   fortran/trans-openmp.cc:omp_map_decayed_kind.  */
+
+static enum gomp_map_kind
+omp_map_decayed_kind (enum gomp_map_kind mapper_kind,
+		      enum gomp_map_kind invoked_as, bool exit_p)
+{
+  if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE)
+    return invoked_as;
+
+  bool force_p, always_p, present_p;
+
+  invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p);
+  gomp_map_kind decay_to;
+
+  switch (mapper_kind)
+    {
+    case GOMP_MAP_ALLOC:
+      if (exit_p && invoked_as == GOMP_MAP_FROM)
+	decay_to = GOMP_MAP_RELEASE;
+      else
+	decay_to = GOMP_MAP_ALLOC;
+      break;
+
+    case GOMP_MAP_TO:
+      if (invoked_as == GOMP_MAP_FROM)
+	decay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+      else if (invoked_as == GOMP_MAP_ALLOC)
+	decay_to = GOMP_MAP_ALLOC;
+      else
+	decay_to = GOMP_MAP_TO;
+      break;
+
+    case GOMP_MAP_FROM:
+      if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO)
+	decay_to = GOMP_MAP_ALLOC;
+      else
+	decay_to = GOMP_MAP_FROM;
+      break;
+
+    case GOMP_MAP_TOFROM:
+    case GOMP_MAP_UNSET:
+      decay_to = invoked_as;
+      break;
+
+    default:
+      gcc_unreachable ();
+    }
+
+  return omp_join_map_kind (decay_to, force_p, always_p, present_p);
+}
+
 /* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
    OUTLIST.  OUTER_KIND is the mapping kind to use if not already specified in
    the mapper declaration.  */
 
 static tree *
 omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
-			enum gomp_map_kind outer_kind)
+			enum gomp_map_kind outer_kind,
+			enum c_omp_region_type ort)
 {
   tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
   tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
@@ -4423,8 +4599,10 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 
       walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
 
-      if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
-	OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
+      enum gomp_map_kind decayed_kind
+	= omp_map_decayed_kind (clause_kind, outer_kind,
+				(ort & C_ORT_EXIT_DATA) != 0);
+      OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind);
 
       type = TYPE_MAIN_VARIANT (type);
 
@@ -4441,11 +4619,8 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 	    = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
 	  if (nested_mapper != mapper)
 	    {
-	      if (clause_kind == GOMP_MAP_UNSET)
-		clause_kind = outer_kind;
-
 	      outlist = omp_instantiate_mapper (outlist, nested_mapper,
-						t, clause_kind);
+						t, outer_kind, ort);
 	      continue;
 	    }
 	}
@@ -4467,7 +4642,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
    visible in the current parsing context.  */
 
 tree
-c_omp_instantiate_mappers (tree clauses)
+c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
 {
   tree c, *pc, mapper_name = NULL_TREE;
 
@@ -4540,7 +4715,7 @@ c_omp_instantiate_mappers (tree clauses)
 	      {
 		tree mapper
 		  = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
-		pc = omp_instantiate_mapper (pc, mapper, t, kind);
+		pc = omp_instantiate_mapper (pc, mapper, t, kind, ort);
 		using_mapper = true;
 	      }
 	    else if (mapper_name)
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index b1301ae..12de2d6 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,15 @@
+2023-07-06  Julian Brown  <julian@codesourcery.com>
+
+	* c-parser.cc (c_parser_omp_target_data): Instantiate mappers for
+	'omp target data'.
+	(c_parser_omp_target_enter_data): Instantiate mappers for 'omp target
+	enter data'.
+	(c_parser_omp_target_exit_data): Instantiate mappers for 'omp target
+	exit data'.
+	(c_parser_omp_target): Add c_omp_region_type argument to
+	c_omp_instantiate_mappers call.
+	* c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype.
+
 2023-06-30  Julian Brown  <julian@codesourcery.com>
 
 	* c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index ebcae16..0cd2cd8 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -24427,7 +24427,9 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
-				"#pragma omp target data");
+				"#pragma omp target data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -24585,7 +24587,9 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
-				"#pragma omp target enter data");
+				"#pragma omp target enter data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -24696,6 +24700,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				"#pragma omp target exit data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);
   clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
@@ -24953,7 +24958,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
-  clauses = c_omp_instantiate_mappers (clauses);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);
   clauses  = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
   c_omp_adjust_map_clauses (clauses, true);
 
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 0c9ce03..d42d259 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -899,7 +899,6 @@ extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);
 extern tree c_omp_mapper_id (tree);
 extern tree c_omp_mapper_decl (tree);
 extern void c_omp_scan_mapper_bindings (location_t, tree *, tree);
-extern tree c_omp_instantiate_mappers (tree);
 extern bool c_check_in_current_scope (tree);
 extern void c_pushtag (location_t, tree, tree);
 extern void c_bind (location_t, tree, bool);
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index ff71a53..c4f4612 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,21 @@
+2023-07-06  Julian Brown  <julian@codesourcery.com>
+
+	* parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp
+	target data'.
+	(cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target
+	enter data'.
+	(cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target
+	exit data'.
+	(cp_parser_omp_target): Add c_omp_region_type argument to
+	c_omp_instantiate_mappers call.
+	* pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other
+	than just C_ORT_OMP_TARGET.
+	(tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE,
+	OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza.
+	* semantics.cc (cxx_omp_map_array_section): Avoid calling
+	build_array_ref for non-array/non-pointer bases (error reported
+	already).
+
 2023-06-30  Julian Brown  <julian@codesourcery.com>
 
 	* constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 217ff03..91a5197 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -47272,7 +47272,10 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
 
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
-				 "#pragma omp target data", pragma_tok);
+				 "#pragma omp target data", pragma_tok, false);
+  if (!processing_template_decl)
+    clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -47387,7 +47390,11 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
 
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
-				 "#pragma omp target enter data", pragma_tok);
+				 "#pragma omp target enter data", pragma_tok,
+				 false);
+  if (!processing_template_decl)
+    clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -47504,6 +47511,8 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				 "#pragma omp target exit data", pragma_tok,
 				 false);
+  if (!processing_template_decl)
+    clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);
   clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
@@ -47798,7 +47807,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
   if (!processing_template_decl)
-    clauses = c_omp_instantiate_mappers (clauses);
+    clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);
   clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
 
   c_omp_adjust_map_clauses (clauses, true);
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 241e816..fa3dc00 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17868,8 +17868,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
   new_clauses = nreverse (new_clauses);
   if (ort != C_ORT_OMP_DECLARE_SIMD)
     {
-      if (ort == C_ORT_OMP_TARGET)
-	new_clauses = c_omp_instantiate_mappers (new_clauses);
+      if (ort & C_ORT_OMP)
+	new_clauses = c_omp_instantiate_mappers (new_clauses, ort);
       new_clauses = finish_omp_clauses (new_clauses, ort);
       if (linear_no_step)
 	for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -19295,7 +19295,9 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)
     case OMP_TARGET_UPDATE:
     case OMP_TARGET_ENTER_DATA:
     case OMP_TARGET_EXIT_DATA:
-      tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), C_ORT_OMP, args,
+      tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t),
+				(TREE_CODE (t) == OMP_TARGET_EXIT_DATA
+				 ? C_ORT_OMP_EXIT_DATA : C_ORT_OMP), args,
 				complain, in_decl);
       t = copy_node (t);
       OMP_STANDALONE_CLAUSES (t) = tmp;
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index dc92a71..65df7aa 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -6454,7 +6454,10 @@ cxx_omp_map_array_section (location_t loc, tree t)
       if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
 	t = convert_from_reference (t);
 
-      t = build_array_ref (loc, t, low);
+      if (TYPE_PTR_P (TREE_TYPE (t)))
+	t = build_array_ref (loc, t, low);
+      else
+	t = error_mark_node;
     }
 
   return t;
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index f9b90ac..ce5eaa1 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,17 @@
+2023-07-06  Julian Brown  <julian@codesourcery.com>
+
+	* trans-openmp.cc (omp_split_map_op, omp_join_map_op,
+	omp_map_decayed_kind): New functions.
+	(gfc_trans_omp_instantiate_mapper): Add CD parameter.  Implement map
+	kind decay.
+	(gfc_trans_omp_instantiate_mappers): Add CD parameter.  Pass to above
+	function.
+	(gfc_trans_omp_target_data): Instantiate mappers for 'omp target data'.
+	(gfc_trans_omp_target_enter_data): Instantiate mappers for 'omp target
+	enter data'.
+	(gfc_trans_omp_target_exit_data): Instantiate mappers for 'omp target
+	exit data'.
+
 2023-06-30  Julian Brown  <julian@codesourcery.com>
 
 	* dump-parse-tree.cc (show_attr): Show omp_udm_artificial_var flag.
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 609604f..831b34b 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -9525,6 +9525,180 @@ gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses *clausesa,
   return gfc_finish_block (&block);
 }
 
+static enum gfc_omp_map_op
+omp_split_map_op (enum gfc_omp_map_op op, bool *force_p, bool *always_p,
+		  bool *present_p)
+{
+  *force_p = *always_p = *present_p = false;
+
+  switch (op)
+    {
+    case OMP_MAP_FORCE_ALLOC:
+    case OMP_MAP_FORCE_TO:
+    case OMP_MAP_FORCE_FROM:
+    case OMP_MAP_FORCE_TOFROM:
+    case OMP_MAP_FORCE_PRESENT:
+      *force_p = true;
+      break;
+    case OMP_MAP_ALWAYS_TO:
+    case OMP_MAP_ALWAYS_FROM:
+    case OMP_MAP_ALWAYS_TOFROM:
+      *always_p = true;
+      break;
+    case OMP_MAP_ALWAYS_PRESENT_TO:
+    case OMP_MAP_ALWAYS_PRESENT_FROM:
+    case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+      *always_p = true;
+      /* Fallthrough.  */
+    case OMP_MAP_PRESENT_ALLOC:
+    case OMP_MAP_PRESENT_TO:
+    case OMP_MAP_PRESENT_FROM:
+    case OMP_MAP_PRESENT_TOFROM:
+      *present_p = true;
+      break;
+    default:
+      ;
+    }
+
+  switch (op)
+    {
+    case OMP_MAP_ALLOC:
+    case OMP_MAP_FORCE_ALLOC:
+    case OMP_MAP_PRESENT_ALLOC:
+      return OMP_MAP_ALLOC;
+    case OMP_MAP_TO:
+    case OMP_MAP_FORCE_TO:
+    case OMP_MAP_ALWAYS_TO:
+    case OMP_MAP_PRESENT_TO:
+    case OMP_MAP_ALWAYS_PRESENT_TO:
+      return OMP_MAP_TO;
+    case OMP_MAP_FROM:
+    case OMP_MAP_FORCE_FROM:
+    case OMP_MAP_ALWAYS_FROM:
+    case OMP_MAP_PRESENT_FROM:
+    case OMP_MAP_ALWAYS_PRESENT_FROM:
+      return OMP_MAP_FROM;
+    case OMP_MAP_TOFROM:
+    case OMP_MAP_FORCE_TOFROM:
+    case OMP_MAP_ALWAYS_TOFROM:
+    case OMP_MAP_PRESENT_TOFROM:
+    case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+      return OMP_MAP_TOFROM;
+    default:
+      ;
+    }
+  return op;
+}
+
+static enum gfc_omp_map_op
+omp_join_map_op (enum gfc_omp_map_op op, bool force_p, bool always_p,
+		 bool present_p)
+{
+  gcc_assert (!force_p || !(always_p || present_p));
+
+  switch (op)
+    {
+    case OMP_MAP_ALLOC:
+      if (force_p)
+	return OMP_MAP_FORCE_ALLOC;
+      else if (present_p)
+	return OMP_MAP_PRESENT_ALLOC;
+      break;
+
+    case OMP_MAP_TO:
+      if (force_p)
+	return OMP_MAP_FORCE_TO;
+      else if (always_p && present_p)
+	return OMP_MAP_ALWAYS_PRESENT_TO;
+      else if (always_p)
+	return OMP_MAP_ALWAYS_TO;
+      else if (present_p)
+	return OMP_MAP_PRESENT_TO;
+      break;
+
+    case OMP_MAP_FROM:
+      if (force_p)
+	return OMP_MAP_FORCE_FROM;
+      else if (always_p && present_p)
+	return OMP_MAP_ALWAYS_PRESENT_FROM;
+      else if (always_p)
+	return OMP_MAP_ALWAYS_FROM;
+      else if (present_p)
+	return OMP_MAP_PRESENT_FROM;
+      break;
+
+    case OMP_MAP_TOFROM:
+      if (force_p)
+	return OMP_MAP_FORCE_TOFROM;
+      else if (always_p && present_p)
+	return OMP_MAP_ALWAYS_PRESENT_TOFROM;
+      else if (always_p)
+	return OMP_MAP_ALWAYS_TOFROM;
+      else if (present_p)
+	return OMP_MAP_PRESENT_TOFROM;
+      break;
+
+    default:
+      ;
+    }
+
+  return op;
+}
+
+/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive").  Return the
+   map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS
+   specified on the clause that invokes the mapper.  See also
+   c-family/c-omp.cc:omp_map_decayed_kind.  */
+
+static enum gfc_omp_map_op
+omp_map_decayed_kind (enum gfc_omp_map_op mapper_kind,
+		      enum gfc_omp_map_op invoked_as, bool exit_p)
+{
+  if (invoked_as == OMP_MAP_RELEASE || invoked_as == OMP_MAP_DELETE)
+    return invoked_as;
+
+  bool force_p, always_p, present_p;
+
+  invoked_as = omp_split_map_op (invoked_as, &force_p, &always_p, &present_p);
+  gfc_omp_map_op decay_to;
+
+  switch (mapper_kind)
+    {
+    case OMP_MAP_ALLOC:
+      if (exit_p && invoked_as == OMP_MAP_FROM)
+	decay_to = OMP_MAP_RELEASE;
+      else
+	decay_to = OMP_MAP_ALLOC;
+      break;
+
+    case OMP_MAP_TO:
+      if (invoked_as == OMP_MAP_FROM)
+	decay_to = exit_p ? OMP_MAP_RELEASE : OMP_MAP_ALLOC;
+      else if (invoked_as == OMP_MAP_ALLOC)
+	decay_to = OMP_MAP_ALLOC;
+      else
+	decay_to = OMP_MAP_TO;
+      break;
+
+    case OMP_MAP_FROM:
+      if (invoked_as == OMP_MAP_ALLOC || invoked_as == OMP_MAP_TO)
+	decay_to = OMP_MAP_ALLOC;
+      else
+	decay_to = OMP_MAP_FROM;
+      break;
+
+    case OMP_MAP_TOFROM:
+    case OMP_MAP_UNSET:
+      decay_to = invoked_as;
+      break;
+
+    default:
+      gcc_unreachable ();
+    }
+
+  return omp_join_map_op (decay_to, force_p, always_p, present_p);
+}
+
 static gfc_symtree *gfc_subst_replace;
 static gfc_ref *gfc_subst_prepend_ref;
 
@@ -9591,7 +9765,8 @@ gfc_subst_mapper_var (gfc_symbol **out_sym, gfc_expr **out_expr,
 
 static gfc_omp_namelist **
 gfc_trans_omp_instantiate_mapper (gfc_omp_namelist **outlistp,
-				  gfc_omp_namelist *clause, gfc_omp_udm *udm)
+				  gfc_omp_namelist *clause, gfc_omp_udm *udm,
+				  toc_directive cd)
 {
   /* Here "sym" and "expr" describe the clause as written, to be substituted
      for the dummy variable in the mapper definition.  */
@@ -9672,10 +9847,10 @@ gfc_trans_omp_instantiate_mapper (gfc_omp_namelist **outlistp,
 			    sym, expr, udm->var_sym, mapper_clause->sym,
 			    mapper_clause->expr);
 
-      if (mapper_clause->u.map_op == OMP_MAP_UNSET)
-	new_clause->u.map_op = outer_map_op;
-      else
-	new_clause->u.map_op = mapper_clause->u.map_op;
+      enum gfc_omp_map_op map_clause_op = mapper_clause->u.map_op;
+      new_clause->u.map_op
+	= omp_map_decayed_kind (map_clause_op, outer_map_op,
+				(cd == TOC_OPENMP_EXIT_DATA));
 
       new_clause->where = clause->where;
 
@@ -9684,7 +9859,7 @@ gfc_trans_omp_instantiate_mapper (gfc_omp_namelist **outlistp,
 	{
 	  gfc_omp_udm *inner_udm = mapper_clause->u2.udm->udm;
 	  outlistp = gfc_trans_omp_instantiate_mapper (outlistp, new_clause,
-						       inner_udm);
+						       inner_udm, cd);
 	}
       else
 	{
@@ -9697,7 +9872,8 @@ gfc_trans_omp_instantiate_mapper (gfc_omp_namelist **outlistp,
 }
 
 static void
-gfc_trans_omp_instantiate_mappers (gfc_omp_clauses *clauses)
+gfc_trans_omp_instantiate_mappers (gfc_omp_clauses *clauses,
+				   toc_directive cd = TOC_OPENMP)
 {
   gfc_omp_namelist *clause = clauses->lists[OMP_LIST_MAP];
   gfc_omp_namelist **clausep = &clauses->lists[OMP_LIST_MAP];
@@ -9706,9 +9882,8 @@ gfc_trans_omp_instantiate_mappers (gfc_omp_clauses *clauses)
     {
       if (clause->u2.udm)
 	{
-	  clausep = gfc_trans_omp_instantiate_mapper (clausep,
-						      clause,
-						      clause->u2.udm->udm);
+	  clausep = gfc_trans_omp_instantiate_mapper (clausep, clause,
+						      clause->u2.udm->udm, cd);
 	  *clausep = clause->next;
 	}
       else
@@ -10160,8 +10335,9 @@ gfc_trans_omp_target_data (gfc_code *code)
   tree stmt, omp_clauses;
 
   gfc_start_block (&block);
-  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-				       code->loc);
+  gfc_omp_clauses *target_data_clauses = code->ext.omp_clauses;
+  gfc_trans_omp_instantiate_mappers (target_data_clauses);
+  omp_clauses = gfc_trans_omp_clauses (&block, target_data_clauses, code->loc);
   stmt = gfc_trans_omp_code (code->block->next, true);
   stmt = build2_loc (gfc_get_location (&code->loc), OMP_TARGET_DATA,
 		     void_type_node, stmt, omp_clauses);
@@ -10176,7 +10352,9 @@ gfc_trans_omp_target_enter_data (gfc_code *code)
   tree stmt, omp_clauses;
 
   gfc_start_block (&block);
-  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+  gfc_omp_clauses *target_enter_data_clauses = code->ext.omp_clauses;
+  gfc_trans_omp_instantiate_mappers (target_enter_data_clauses);
+  omp_clauses = gfc_trans_omp_clauses (&block, target_enter_data_clauses,
 				       code->loc);
   stmt = build1_loc (input_location, OMP_TARGET_ENTER_DATA, void_type_node,
 		     omp_clauses);
@@ -10191,7 +10369,10 @@ gfc_trans_omp_target_exit_data (gfc_code *code)
   tree stmt, omp_clauses;
 
   gfc_start_block (&block);
-  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+  gfc_omp_clauses *target_exit_data_clauses = code->ext.omp_clauses;
+  gfc_trans_omp_instantiate_mappers (target_exit_data_clauses,
+				     TOC_OPENMP_EXIT_DATA);
+  omp_clauses = gfc_trans_omp_clauses (&block, target_exit_data_clauses,
 				       code->loc, TOC_OPENMP_EXIT_DATA);
   stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node,
 		     omp_clauses);
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 7359347..2c1e857 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,11 @@
+2023-07-06  Julian Brown  <julian@codesourcery.com>
+
+	* c-c++-common/gomp/declare-mapper-15.c: New test.
+	* c-c++-common/gomp/declare-mapper-16.c: New test.
+	* g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output.
+	* gfortran.dg/gomp/declare-mapper-22.f90: New test.
+	* gfortran.dg/gomp/declare-mapper-23.f90: New test.
+
 2023-06-30  Julian Brown  <julian@codesourcery.com>
 
 	* gfortran.dg/gomp/declare-mapper-1.f90: New test.
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
new file mode 100644
index 0000000..ecda2e5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
@@ -0,0 +1,59 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+typedef struct {
+  int a, b, c, d;
+} S;
+
+int main ()
+{
+  S s;
+  #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \
+				   map(from: x.c) map(tofrom: x.d)
+
+  #pragma omp target enter data map(to: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  #pragma omp target exit data map(from: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: 4\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+  #pragma omp target enter data map(alloc: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  #pragma omp target exit data map(release: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(release:s\.c \[len: [0-9]+\]\) map\(release:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+  #pragma omp target enter data map(present, to: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(force_present:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  #pragma omp target exit data map(present, from: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+  #pragma omp target enter data map(always, to: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(always,to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(always,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  #pragma omp target exit data map(always, from: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,from:s\.c \[len: [0-9]+\]\) map\(always,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+  #pragma omp target enter data map(always, present, to: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(always,present,to:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(always,present,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  #pragma omp target exit data map(always, present, from: s)
+
+  /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,present,from:s\.c \[len: [0-9]+\]\) map\(always,present,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c
new file mode 100644
index 0000000..20383cc
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c
@@ -0,0 +1,39 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+typedef struct {
+  int a, b, c, d;
+} S;
+
+int main ()
+{
+  S s = { 0, 0, 0, 0 };
+  #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \
+				   map(from: x.c) map(tofrom: x.d)
+
+  #pragma omp target data map(s)
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(tofrom:s\.d \[len: [0-9]+\]\)} 3 "gimple" } } */
+  {
+    #pragma omp target
+    {
+      s.a++;
+      s.b++;
+      s.c++;
+      s.d++;
+    }
+  }
+
+  #pragma omp target data map(alloc: s)
+  /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+  {
+    #pragma omp target
+    {
+      s.a++;
+      s.b++;
+      s.c++;
+      s.d++;
+    }
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
index 3177d20..8af3bac 100644
--- a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
@@ -55,4 +55,4 @@ int main (int argc, char *argv[])
 }
 
 // { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } }
-// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(alloc:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-22.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-22.f90
new file mode 100644
index 0000000..483e848
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-22.f90
@@ -0,0 +1,60 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+type t
+integer, allocatable :: arrcomp(:)
+integer :: b, c, d
+end type t
+
+type(t) :: myvar
+
+!$omp declare mapper (t :: x) map(to: x%arrcomp) map(alloc: x%b) &
+!$omp &                       map(from: x%c) map(tofrom: x%d)
+
+allocate (myvar%arrcomp(1:100))
+
+!$omp target enter data map(to: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(to:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(alloc:myvar\.b \[len: [0-9]+\]\) map\(alloc:myvar\.c \[len: 4\]\) map\(to:myvar\.d \[len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+!$omp target exit data map(from: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(release:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(release:myvar\.b \[len: [0-9]+\]\) map\(from:myvar\.c \[len: [0-9]+\]\) map\(from:myvar\.d \[len: [0-9]+\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+
+!$omp target enter data map(alloc: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(to:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(alloc:myvar\.b \[len: [0-9]+\]\) map\(alloc:myvar\.c \[len: [0-9]+\]\) map\(alloc:myvar\.d \[len: [0-9]+\]\) map\(alloc:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+!$omp target exit data map(release: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(release:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(release:myvar\.b \[len: [0-9]+\]\) map\(release:myvar\.c \[len: [0-9]+\]\) map\(release:myvar\.d \[len: [0-9]+\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+
+!$omp target enter data map(present, to: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(force_present:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:myvar\.arrcomp\.data \[bias: 0\]\) map\(struct:myvar \[len: 4\]\) map\(to:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(force_present:myvar\.b \[len: [0-9]+\]\) map\(force_present:myvar\.c \[len: [0-9]+\]\) map\(force_present:myvar\.d \[len: [0-9]+\]\)} 1 "gimple" } }
+
+!$omp target exit data map(present, from: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(release:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(release:myvar\.b \[len: [0-9]+\]\) map\(force_present:myvar\.c \[len: [0-9]+\]\) map\(force_present:myvar\.d \[len: [0-9]+\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+
+!$omp target enter data map(always, to: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(to:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(alloc:myvar\.b \[len: [0-9]+\]\) map\(alloc:myvar\.c \[len: [0-9]+\]\) map\(always,to:myvar\.d \[len: [0-9]+\]\) map\(always,to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+!$omp target exit data map(always, from: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(release:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(release:myvar\.b \[len: [0-9]+\]\) map\(always,from:myvar\.c \[len: [0-9]+\]\) map\(always,from:myvar\.d \[len: [0-9]+\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+
+!$omp target enter data map(always, present, to: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(to:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(force_present:myvar\.b \[len: [0-9]+\]\) map\(force_present:myvar\.c \[len: [0-9]+\]\) map\(always,present,to:myvar\.d \[len: [0-9]+\]\) map\(always,present,to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+!$omp target exit data map(always, present, from: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(release:myvar\.arrcomp \[pointer set, len: [0-9]+\]\) map\(release:myvar\.b \[len: [0-9]+\]\) map\(always,present,from:myvar\.c \[len: [0-9]+\]\) map\(always,present,from:myvar\.d \[len: [0-9]+\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\)} 1 "gimple" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-23.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-23.f90
new file mode 100644
index 0000000..6c07261
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-23.f90
@@ -0,0 +1,25 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+type t
+integer :: a, b, c, d
+end type t
+
+type(t) :: myvar
+
+!$omp declare mapper (t :: x) map(to: x%a) map(alloc: x%b) &
+!$omp &                       map(from: x%c) map(tofrom: x%d)
+
+!$omp target data map(to: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(to:myvar\.a \[len: [0-9]+\]\) map\(alloc:myvar\.b \[len: [0-9]+\]\) map\(alloc:myvar\.c \[len: [0-9]+\]\) map\(to:myvar\.d \[len: [0-9]+\]\)} 1 "gimple" } }
+
+!$omp end target data
+
+!$omp target data map(alloc: myvar)
+
+! { dg-final { scan-tree-dump-times {map\(struct:myvar \[len: 4\]\) map\(alloc:myvar\.a \[len: [0-9]+\]\) map\(alloc:myvar\.b \[len: [0-9]+\]\) map\(alloc:myvar\.c \[len: [0-9]+\]\) map\(alloc:myvar\.d \[len: [0-9]+\]\)} 1 "gimple" } }
+
+!$omp end target data
+
+end