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

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index ea6c479cd62..c805c8b2f7e 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1270,8 +1270,10 @@  enum c_omp_region_type
   C_ORT_ACC			= 1 << 1,
   C_ORT_DECLARE_SIMD		= 1 << 2,
   C_ORT_TARGET			= 1 << 3,
+  C_ORT_EXIT_DATA		= 1 << 4,
   C_ORT_OMP_DECLARE_SIMD	= C_ORT_OMP | C_ORT_DECLARE_SIMD,
   C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET,
+  C_ORT_OMP_EXIT_DATA		= C_ORT_OMP | C_ORT_EXIT_DATA,
   C_ORT_ACC_TARGET		= C_ORT_ACC | C_ORT_TARGET
 };
 
@@ -1310,7 +1312,7 @@  extern void c_oacc_annotate_loops_in_kernels_regions (tree, 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 9e307411f76..2ddb30b54f8 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -5001,13 +5001,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);
@@ -5064,8 +5240,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);
 
@@ -5082,11 +5260,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;
 	    }
 	}
@@ -5108,7 +5283,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;
 
@@ -5181,7 +5356,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/c-parser.cc b/gcc/c/c-parser.cc
index 7e895e11da2..f1348a7e5f3 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -22931,7 +22931,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;)
@@ -23118,7 +23120,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;)
@@ -23228,7 +23232,9 @@  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");
+				"#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;
   for (tree *pc = &clauses; *pc;)
@@ -23485,7 +23491,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 ee29f9de2cc..378707e5155 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -888,7 +888,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/parser.cc b/gcc/cp/parser.cc
index 92495c1049a..4984accb0bd 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -46598,7 +46598,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;)
@@ -46712,7 +46715,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;)
@@ -46827,7 +46834,11 @@  cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
 
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
-				 "#pragma omp target exit data", pragma_tok);
+				 "#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;
   for (tree *pc = &clauses; *pc;)
@@ -47151,7 +47162,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 83c25811222..ba78955ccdd 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -18319,8 +18319,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))
@@ -19736,7 +19736,9 @@  tsubst_expr (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 e55ff6d5a2d..e4feb0cc468 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -6261,7 +6261,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/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index d30130bc140..6cb5340687e 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -10084,6 +10084,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;
 
@@ -10150,7 +10324,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.  */
@@ -10231,10 +10406,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;
 
@@ -10243,7 +10418,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
 	{
@@ -10256,7 +10431,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];
@@ -10265,9 +10441,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
@@ -10721,8 +10896,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);
@@ -10737,7 +10913,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);
@@ -10752,7 +10930,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/c-c++-common/gomp/declare-mapper-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
new file mode 100644
index 00000000000..ecda2e5ebd1
--- /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 00000000000..20383cc2d69
--- /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 3177d20adbc..8af3bac1071 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 00000000000..0564e564395
--- /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\.b \[len: [0-9]+\]\) map\(from:myvar\.c \[len: [0-9]+\]\) map\(from:myvar\.d \[len: [0-9]+\]\) map\(release:myvar\.arrcomp \[len: [0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\)} 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\.b \[len: [0-9]+\]\) map\(release:myvar\.c \[len: [0-9]+\]\) map\(release:myvar\.d \[len: [0-9]+\]\) map\(release:myvar\.arrcomp \[len: [0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\)} 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\.b \[len: [0-9]+\]\) map\(force_present:myvar\.c \[len: [0-9]+\]\) map\(force_present:myvar\.d \[len: [0-9]+\]\) map\(release:myvar\.arrcomp \[len: [0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\)} 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\.b \[len: [0-9]+\]\) map\(always,from:myvar\.c \[len: [0-9]+\]\) map\(always,from:myvar\.d \[len: [0-9]+\]\) map\(release:myvar\.arrcomp \[len: [0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\)} 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\.b \[len: [0-9]+\]\) map\(always,present,from:myvar\.c \[len: [0-9]+\]\) map\(always,present,from:myvar\.d \[len: [0-9]+\]\) map\(release:myvar\.arrcomp \[len: [0-9]+\]\) map\(detach:myvar\.arrcomp\.data \[bias: 0\]\) map\(release:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\)} 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 00000000000..6c07261040d
--- /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