diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc
index c05be833357..050049f3c1a 100644
--- a/gcc/cp/cp-gimplify.cc
+++ b/gcc/cp/cp-gimplify.cc
@@ -2282,6 +2282,12 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */)
     }
 }
 
+tree
+cxx_omp_finish_mapper_clauses (tree clauses)
+{
+  return finish_omp_clauses (clauses, C_ORT_OMP);
+}
+
 /* Return true if DECL's DECL_VALUE_EXPR (if any) should be
    disregarded in OpenMP construct, because it is going to be
    remapped during OpenMP lowering.  SHARED is true if DECL
diff --git a/gcc/cp/cp-objcp-common.h b/gcc/cp/cp-objcp-common.h
index 1a67f14d9b3..b6f72e004d3 100644
--- a/gcc/cp/cp-objcp-common.h
+++ b/gcc/cp/cp-objcp-common.h
@@ -185,6 +185,8 @@ extern tree cxx_simulate_record_decl (location_t, const char *,
 #define LANG_HOOKS_OMP_CLAUSE_DTOR cxx_omp_clause_dtor
 #undef LANG_HOOKS_OMP_FINISH_CLAUSE
 #define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
+#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES cxx_omp_finish_mapper_clauses
 #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
 #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
 #undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index c897da204fe..be41019529d 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -2893,6 +2893,7 @@ struct GTY(()) lang_decl_fn {
 
   unsigned this_thunk_p : 1;
   unsigned omp_declare_reduction_p : 1;
+  unsigned omp_declare_mapper_p : 1;
   unsigned has_dependent_explicit_spec_p : 1;
   unsigned immediate_fn_p : 1;
   unsigned maybe_deleted : 1;
@@ -4289,6 +4290,11 @@ get_vec_init_expr (tree t)
 #define DECL_OMP_DECLARE_REDUCTION_P(NODE) \
   (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p)
 
+/* Nonzero if NODE is an artificial FUNCTION_DECL for
+   #pragma omp declare mapper.  */
+#define DECL_OMP_DECLARE_MAPPER_P(NODE) \
+  (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_mapper_p)
+
 /* Nonzero if DECL has been declared threadprivate by
    #pragma omp threadprivate.  */
 #define CP_DECL_THREADPRIVATE_P(DECL) \
@@ -7657,10 +7663,13 @@ extern tree finish_qualified_id_expr		(tree, tree, bool, bool,
 extern void simplify_aggr_init_expr		(tree *);
 extern void finalize_nrv			(tree *, tree, tree);
 extern tree omp_reduction_id			(enum tree_code, tree, tree);
+extern tree omp_mapper_id			(tree, tree);
 extern tree cp_remove_omp_priv_cleanup_stmt	(tree *, int *, void *);
 extern bool cp_check_omp_declare_reduction	(tree);
+extern bool cp_check_omp_declare_mapper		(tree);
 extern void finish_omp_declare_simd_methods	(tree);
 extern tree finish_omp_clauses			(tree, enum c_omp_region_type);
+extern tree omp_instantiate_mappers		(tree);
 extern tree push_omp_privatization_clauses	(bool);
 extern void pop_omp_privatization_clauses	(tree);
 extern void save_omp_privatization_clauses	(vec<tree> &);
@@ -8212,6 +8221,7 @@ extern tree cxx_omp_clause_copy_ctor		(tree, tree, tree);
 extern tree cxx_omp_clause_assign_op		(tree, tree, tree);
 extern tree cxx_omp_clause_dtor			(tree, tree);
 extern void cxx_omp_finish_clause		(tree, gimple_seq *, bool);
+extern tree cxx_omp_finish_mapper_clauses	(tree);
 extern bool cxx_omp_privatize_by_reference	(const_tree);
 extern bool cxx_omp_disregard_value_expr	(tree, bool);
 extern void cp_fold_function			(tree);
diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc
index b72b2a8456b..48ca679a91a 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -1897,6 +1897,18 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
 	      "previous %<pragma omp declare reduction%> declaration");
       return error_mark_node;
     }
+  else if (TREE_CODE (newdecl) == FUNCTION_DECL
+	   && DECL_OMP_DECLARE_MAPPER_P (newdecl))
+    {
+      /* OMP UDMs are never duplicates either.  */
+      gcc_assert (DECL_OMP_DECLARE_MAPPER_P (olddecl));
+      error_at (newdecl_loc,
+		"redeclaration of %<pragma omp declare mapper%>");
+      inform (olddecl_loc,
+	      "previous %<pragma omp declare mapper%> declaration");
+      return error_mark_node;
+
+    }
   else if (TREE_CODE (newdecl) == FUNCTION_DECL
 	    && ((DECL_TEMPLATE_SPECIALIZATION (olddecl)
 		 && (!DECL_TEMPLATE_INFO (newdecl)
@@ -17977,7 +17989,8 @@ finish_function (bool inline_p)
   /* Perform delayed folding before NRV transformation.  */
   if (!processing_template_decl
       && !DECL_IMMEDIATE_FUNCTION_P (fndecl)
-      && !DECL_OMP_DECLARE_REDUCTION_P (fndecl))
+      && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)
+      && !DECL_OMP_DECLARE_MAPPER_P (fndecl))
     cp_fold_function (fndecl);
 
   /* Set up the named return value optimization, if we can.  Candidate
@@ -18054,7 +18067,8 @@ finish_function (bool inline_p)
   /* Genericize before inlining.  */
   if (!processing_template_decl
       && !DECL_IMMEDIATE_FUNCTION_P (fndecl)
-      && !DECL_OMP_DECLARE_REDUCTION_P (fndecl))
+      && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)
+      && !DECL_OMP_DECLARE_MAPPER_P (fndecl))
     cp_genericize (fndecl);
 
   /* Emit the resumer and destroyer functions now, providing that we have
diff --git a/gcc/cp/mangle.cc b/gcc/cp/mangle.cc
index 75388e99bfd..d06bff22f23 100644
--- a/gcc/cp/mangle.cc
+++ b/gcc/cp/mangle.cc
@@ -955,10 +955,11 @@ decl_mangling_context (tree decl)
 
   tcontext = CP_DECL_CONTEXT (decl);
 
-  /* Ignore the artificial declare reduction functions.  */
+  /* Ignore the artificial declare reduction and declare mapper functions.  */
   if (tcontext
       && TREE_CODE (tcontext) == FUNCTION_DECL
-      && DECL_OMP_DECLARE_REDUCTION_P (tcontext))
+      && (DECL_OMP_DECLARE_REDUCTION_P (tcontext)
+	  || DECL_OMP_DECLARE_MAPPER_P (tcontext)))
     return decl_mangling_context (tcontext);
 
   return tcontext;
diff --git a/gcc/cp/name-lookup.cc b/gcc/cp/name-lookup.cc
index ce622761a1a..926bf10cc2f 100644
--- a/gcc/cp/name-lookup.cc
+++ b/gcc/cp/name-lookup.cc
@@ -3343,7 +3343,8 @@ set_decl_context_in_fn (tree ctx, tree decl)
     gcc_checking_assert (DECL_LOCAL_DECL_P (decl)
 			 && (DECL_NAMESPACE_SCOPE_P (decl)
 			     || (TREE_CODE (decl) == FUNCTION_DECL
-				 && DECL_OMP_DECLARE_REDUCTION_P (decl))));
+				 && (DECL_OMP_DECLARE_REDUCTION_P (decl)
+				     || DECL_OMP_DECLARE_MAPPER_P (decl)))));
 
   if (!DECL_CONTEXT (decl)
       /* When parsing the parameter list of a function declarator,
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 80d33c3b674..e87de0a38a6 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -26407,10 +26407,12 @@ cp_parser_class_specifier (cp_parser* parser)
 	{
 	  /* OpenMP UDRs need to be parsed before all other functions.  */
 	  FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl)
-	    if (DECL_OMP_DECLARE_REDUCTION_P (decl))
+	    if (DECL_OMP_DECLARE_REDUCTION_P (decl)
+		|| DECL_OMP_DECLARE_MAPPER_P (decl))
 	      cp_parser_late_parsing_for_member (parser, decl);
 	  FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl)
-	    if (!DECL_OMP_DECLARE_REDUCTION_P (decl))
+	    if (!DECL_OMP_DECLARE_REDUCTION_P (decl)
+		&& !DECL_OMP_DECLARE_MAPPER_P (decl))
 	      cp_parser_late_parsing_for_member (parser, decl);
 	}
       else
@@ -32392,6 +32394,8 @@ cp_parser_enclosed_template_argument_list (cp_parser* parser)
   return arguments;
 }
 
+static bool cp_parser_omp_declare_mapper_maplist (tree, cp_parser *);
+
 /* MEMBER_FUNCTION is a member function, or a friend.  If default
    arguments, or the body of the function have not yet been parsed,
    parse them now.  */
@@ -32454,6 +32458,16 @@ cp_parser_late_parsing_for_member (cp_parser* parser, tree member_function)
 	  finish_function (/*inline_p=*/true);
 	  cp_check_omp_declare_reduction (member_function);
 	}
+      else if (DECL_OMP_DECLARE_MAPPER_P (member_function))
+	{
+	  parser->lexer->in_pragma = true;
+	  cp_parser_omp_declare_mapper_maplist (member_function, parser);
+	  finish_function (/*inline_p=*/true);
+	  cp_check_omp_declare_mapper (member_function);
+	  /* If this is a template class, this forces the body of the mapper
+	     to be instantiated.  */
+	  DECL_PRESERVE_P (member_function) = 1;
+	}
       else
 	/* Now, parse the body of the function.  */
 	cp_parser_function_definition_after_declarator (parser,
@@ -39887,13 +39901,12 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
    map ( [map-type-modifier[,] ...] map-kind: variable-list )
 
    map-type-modifier:
-     always | close */
+     always | close | mapper ( mapper-name )  */
 
 static tree
-cp_parser_omp_clause_map (cp_parser *parser, tree list)
+cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind)
 {
   tree nlist, c;
-  enum gomp_map_kind kind = GOMP_MAP_TOFROM;
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
@@ -39911,11 +39924,27 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 
       if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
 	pos++;
+      else if ((cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
+		== CPP_OPEN_PAREN)
+	       && ((cp_lexer_peek_nth_token (parser->lexer, pos + 2)->type
+		    == CPP_NAME)
+		   || ((cp_lexer_peek_nth_token (parser->lexer, pos + 2)->type
+			== CPP_KEYWORD)
+		       && (cp_lexer_peek_nth_token (parser->lexer,
+						    pos + 2)->keyword
+			   == RID_DEFAULT)))
+	       && (cp_lexer_peek_nth_token (parser->lexer, pos + 3)->type
+		   == CPP_CLOSE_PAREN)
+	       && (cp_lexer_peek_nth_token (parser->lexer, pos + 4)->type
+		   == CPP_COMMA))
+	pos += 4;
       pos++;
     }
 
   bool always_modifier = false;
   bool close_modifier = false;
+  bool mapper_modifier = false;
+  tree mapper_name = NULL_TREE;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -39938,6 +39967,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 	      return list;
 	    }
 	  always_modifier = true;
+	  cp_lexer_consume_token (parser->lexer);
 	}
       else if (strcmp ("close", p) == 0)
 	{
@@ -39951,20 +39981,83 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 	      return list;
 	    }
 	  close_modifier = true;
+	  cp_lexer_consume_token (parser->lexer);
+	}
+      else if (strcmp ("mapper", p) == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+
+	  matching_parens parens;
+	  if (parens.require_open (parser))
+	    {
+	      if (mapper_modifier)
+		{
+		  cp_parser_error (parser, "too many %<mapper%> modifiers");
+		  /* Assume it's a well-formed mapper modifier, even if it
+		     seems to be in the wrong place.  */
+		  cp_lexer_consume_token (parser->lexer);
+		  parens.require_close (parser);
+		  cp_parser_skip_to_closing_parenthesis (parser,
+							 /*recovering=*/true,
+							 /*or_comma=*/false,
+							 /*consume_paren=*/
+							 true);
+		  return list;
+		}
+
+	      tok = cp_lexer_peek_token (parser->lexer);
+	      switch (tok->type)
+		{
+		case CPP_NAME:
+		  {
+		    cp_expr e = cp_parser_identifier (parser);
+		    if (e != error_mark_node)
+		      mapper_name = e;
+		    else
+		      goto err;
+		  }
+		break;
+
+		case CPP_KEYWORD:
+		  if (tok->keyword == RID_DEFAULT)
+		    {
+		      cp_lexer_consume_token (parser->lexer);
+		      break;
+		    }
+		  /* Fallthrough.  */
+
+		default:
+		err:
+		  cp_parser_error (parser,
+				   "expected identifier or %<default%>");
+		  return list;
+		}
+
+	      if (!parens.require_close (parser))
+		{
+		  cp_parser_skip_to_closing_parenthesis (parser,
+							 /*recovering=*/true,
+							 /*or_comma=*/false,
+							 /*consume_paren=*/
+							 true);
+		  return list;
+		}
+
+	      mapper_modifier = true;
+	      pos += 3;
+	    }
 	}
       else
 	{
 	  cp_parser_error (parser, "%<#pragma omp target%> with "
-				   "modifier other than %<always%> or "
-				   "%<close%> on %<map%> clause");
+				   "modifier other than %<always%>, %<close%> "
+				   "or %<mapper%> on %<map%> clause");
 	  cp_parser_skip_to_closing_parenthesis (parser,
 						 /*recovering=*/true,
 						 /*or_comma=*/false,
 						 /*consume_paren=*/true);
 	  return list;
 	}
-
-	cp_lexer_consume_token (parser->lexer);
     }
 
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
@@ -40005,8 +40098,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
   nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
 					  NULL, true);
 
+  tree last_new = NULL_TREE;
+
   for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      last_new = c;
+    }
+
+  if (mapper_name)
+    {
+      tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
+      OMP_CLAUSE_DECL (name) = mapper_name;
+      OMP_CLAUSE_CHAIN (name) = nlist;
+      nlist = name;
+
+      gcc_assert (last_new);
+
+      name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
+      OMP_CLAUSE_DECL (name) = null_pointer_node;
+      OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
+      OMP_CLAUSE_CHAIN (last_new) = name;
+    }
 
   return nlist;
 }
@@ -40817,7 +40932,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "detach";
 	  break;
 	case PRAGMA_OMP_CLAUSE_MAP:
-	  clauses = cp_parser_omp_clause_map (parser, clauses);
+	  clauses = cp_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM);
 	  c_name = "map";
 	  break;
 	case PRAGMA_OMP_CLAUSE_DEVICE:
@@ -45006,6 +45121,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
+  if (!processing_template_decl)
+    clauses = omp_instantiate_mappers (clauses);
   clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
 
   c_omp_adjust_map_clauses (clauses, true);
@@ -46954,6 +47071,252 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
   obstack_free (&declarator_obstack, p);
 }
 
+/* OpenMP 5.0:
+   Parse a variable name and a list of map clauses for "omp declare mapper"
+   directives:
+
+   ... var) [clause[[,] clause] ... ] new-line  */
+
+static bool
+cp_parser_omp_declare_mapper_maplist (tree fndecl, cp_parser *parser)
+{
+  pragma_omp_clause c_kind;
+  tree maplist = NULL_TREE, stmt = NULL_TREE;
+  tree mapper_name = NULL_TREE;
+  tree type = TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (fndecl)));
+  tree id = cp_parser_declarator_id (parser, false);
+
+  if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+    return false;
+
+  gcc_assert (TYPE_REF_P (type));
+  type = TREE_TYPE (type);
+
+  keep_next_level (true);
+  tree block = begin_omp_structured_block ();
+
+  tree var = build_lang_decl (VAR_DECL, id, type);
+  pushdecl (var);
+  cp_finish_decl (var, NULL_TREE, 0, NULL_TREE, 0);
+  DECL_ARTIFICIAL (var) = 1;
+  TREE_USED (var) = 1;
+
+  const char *fnname = IDENTIFIER_POINTER (DECL_NAME (fndecl));
+  if (startswith (fnname, "omp declare mapper "))
+    fnname += sizeof "omp declare mapper " - 1;
+  const char *mapname_end = strchr (fnname, '~');
+  if (mapname_end && mapname_end != fnname)
+    {
+      char *tmp = XALLOCAVEC (char, mapname_end - fnname + 1);
+      strncpy (tmp, fnname, mapname_end - fnname);
+      tmp[mapname_end - fnname] = '\0';
+      mapper_name = get_identifier (tmp);
+    }
+
+  while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+    {
+      c_kind = cp_parser_omp_clause_name (parser);
+      if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+	{
+	  if (c_kind != PRAGMA_OMP_CLAUSE_NONE)
+	    cp_parser_error (parser, "unexpected clause");
+	  finish_omp_structured_block (block);
+	  return false;
+	}
+      maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+      if (maplist == NULL_TREE)
+	{
+	  finish_omp_structured_block (block);
+	  return false;
+	}
+    }
+
+  if (maplist == NULL_TREE)
+    {
+      cp_parser_error (parser, "missing %<map%> clause");
+      finish_omp_structured_block (block);
+      return false;
+    }
+
+  stmt = make_node (OMP_DECLARE_MAPPER);
+  TREE_TYPE (stmt) = void_type_node;
+  OMP_DECLARE_MAPPER_ID (stmt) = mapper_name;
+  OMP_DECLARE_MAPPER_TYPE (stmt) = type;
+  OMP_DECLARE_MAPPER_DECL (stmt) = var;
+  OMP_DECLARE_MAPPER_CLAUSES (stmt) = maplist;
+
+  add_stmt (stmt);
+
+  block = finish_omp_structured_block (block);
+
+  add_stmt (block);
+
+  return true;
+}
+
+/* OpenMP 5.0
+   #pragma omp declare mapper([mapper-identifier:]type var) \
+	       [clause[[,] clause] ... ] new-line  */
+
+static void
+cp_parser_omp_declare_mapper (cp_parser *parser, cp_token *pragma_tok,
+			      enum pragma_context)
+{
+  cp_token *token = NULL;
+  cp_token *first_token = NULL;
+  cp_token_cache *cp = NULL;
+  tree type = NULL_TREE, fndecl = NULL_TREE, block = NULL_TREE;
+  bool block_scope = false;
+  /* Don't create location wrapper nodes within "declare mapper"
+     directives.  */
+  auto_suppress_location_wrappers sentinel;
+  tree mapper_name = NULL_TREE;
+  tree mapper_id, fntype;
+
+  if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    goto fail;
+
+  if (current_function_decl)
+    block_scope = true;
+
+  token = cp_lexer_peek_token (parser->lexer);
+
+  if (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+    {
+      switch (token->type)
+	{
+	case CPP_NAME:
+	  {
+	    cp_expr e = cp_parser_identifier (parser);
+	    if (e != error_mark_node)
+	      mapper_name = e;
+	    else
+	      goto fail;
+	  }
+	  break;
+
+	case CPP_KEYWORD:
+	  if (token->keyword == RID_DEFAULT)
+	    {
+	      mapper_name = NULL_TREE;
+	      cp_lexer_consume_token (parser->lexer);
+	      break;
+	    }
+	  /* Fallthrough.  */
+
+	default:
+	  cp_parser_error (parser, "expected identifier or %<default%>");
+	}
+
+      if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
+	goto fail;
+    }
+
+  {
+    const char *saved_message = parser->type_definition_forbidden_message;
+    parser->type_definition_forbidden_message
+      = G_("types may not be defined within %<declare mapper%>");
+    type_id_in_expr_sentinel s (parser);
+    type = cp_parser_type_id (parser);
+    parser->type_definition_forbidden_message = saved_message;
+  }
+
+  if (dependent_type_p (type))
+    mapper_id = omp_mapper_id (mapper_name, NULL_TREE);
+  else
+    mapper_id = omp_mapper_id (mapper_name, type);
+
+  fntype = build_function_type_list (void_type_node,
+				     cp_build_reference_type (type, false),
+				     NULL_TREE);
+  fndecl = build_lang_decl (FUNCTION_DECL, mapper_id, fntype);
+  DECL_SOURCE_LOCATION (fndecl) = pragma_tok->location;
+  DECL_ARTIFICIAL (fndecl) = 1;
+  DECL_EXTERNAL (fndecl) = 1;
+  DECL_DECLARED_INLINE_P (fndecl) = 1;
+  DECL_IGNORED_P (fndecl) = 1;
+  DECL_OMP_DECLARE_MAPPER_P (fndecl) = 1;
+  SET_DECL_ASSEMBLER_NAME (fndecl, get_identifier ("<udm>"));
+  DECL_ATTRIBUTES (fndecl)
+    = tree_cons (get_identifier ("gnu_inline"), NULL_TREE,
+		 DECL_ATTRIBUTES (fndecl));
+
+  if (block_scope)
+    block = begin_omp_structured_block ();
+
+  first_token = cp_lexer_peek_token (parser->lexer);
+
+  if (processing_template_decl)
+    fndecl = push_template_decl (fndecl);
+
+  if (block_scope)
+    {
+      DECL_CONTEXT (fndecl) = current_function_decl;
+      DECL_LOCAL_DECL_P (fndecl) = 1;
+    }
+  else if (current_class_type)
+    {
+      while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+	cp_lexer_consume_token (parser->lexer);
+      cp = cp_token_cache_new (first_token,
+			       cp_lexer_peek_nth_token (parser->lexer, 2));
+      DECL_STATIC_FUNCTION_P (fndecl) = 1;
+      finish_member_declaration (fndecl);
+      DECL_PENDING_INLINE_INFO (fndecl) = cp;
+      DECL_PENDING_INLINE_P (fndecl) = 1;
+      vec_safe_push (unparsed_funs_with_definitions, fndecl);
+      cp_parser_require_pragma_eol (parser, pragma_tok);
+      return;
+    }
+  else
+    {
+      DECL_CONTEXT (fndecl) = current_namespace;
+      tree d = pushdecl (fndecl);
+      gcc_checking_assert (d == error_mark_node || d == fndecl);
+
+      start_preparsed_function (fndecl, NULL_TREE, SF_PRE_PARSED);
+    }
+
+  if (!cp_parser_omp_declare_mapper_maplist (fndecl, parser))
+    {
+      if (block_scope)
+	finish_omp_structured_block (block);
+      else
+	finish_function (false);
+      goto fail;
+    }
+
+  if (!block_scope)
+    {
+      tree fn = finish_function (/*inline_p=*/false);
+      expand_or_defer_fn (fn);
+    }
+  else
+    {
+      DECL_CONTEXT (fndecl) = current_function_decl;
+      if (DECL_TEMPLATE_INFO (fndecl))
+	DECL_CONTEXT (DECL_TI_TEMPLATE (fndecl)) = current_function_decl;
+
+      block = finish_omp_structured_block (block);
+      if (TREE_CODE (block) == BIND_EXPR)
+	DECL_SAVED_TREE (fndecl) = BIND_EXPR_BODY (block);
+      else if (TREE_CODE (block) == STATEMENT_LIST)
+	DECL_SAVED_TREE (fndecl) = block;
+      if (processing_template_decl)
+	add_decl_expr (fndecl);
+      else
+	pushdecl (fndecl);
+    }
+
+  cp_check_omp_declare_mapper (fndecl);
+
+  cp_parser_require_pragma_eol (parser, pragma_tok);
+  return;
+
+fail:
+  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+}
+
 /* OpenMP 4.0
    #pragma omp declare simd declare-simd-clauses[optseq] new-line
    #pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -46994,6 +47357,12 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
 					   context);
 	  return false;
 	}
+      if (strcmp (p, "mapper") == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_parser_omp_declare_mapper (parser, pragma_tok, context);
+	  return false;
+	}
       if (!flag_openmp)  /* flag_openmp_simd  */
 	{
 	  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
@@ -47007,7 +47376,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
 	}
     }
   cp_parser_error (parser, "expected %<simd%>, %<reduction%>, "
-			   "%<target%> or %<variant%>");
+			   "%<target%>, %<mapper%> or %<variant%>");
   cp_parser_require_pragma_eol (parser, pragma_tok);
   return false;
 }
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index a920d86b354..ca7547ec9b0 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -12256,9 +12256,13 @@ instantiate_class_template (tree type)
 	      /* Instantiate members marked with attribute used.  */
 	      if (r != error_mark_node && DECL_PRESERVE_P (r))
 		used.safe_push (r);
-	      if (TREE_CODE (r) == FUNCTION_DECL
-		  && DECL_OMP_DECLARE_REDUCTION_P (r))
-		cp_check_omp_declare_reduction (r);
+	      if (TREE_CODE (r) == FUNCTION_DECL)
+		{
+		  if (DECL_OMP_DECLARE_REDUCTION_P (r))
+		    cp_check_omp_declare_reduction (r);
+		  else if (DECL_OMP_DECLARE_MAPPER_P (r))
+		    cp_check_omp_declare_mapper (r);
+		}
 	    }
 	  else if ((DECL_CLASS_TEMPLATE_P (t) || DECL_IMPLICIT_TYPEDEF_P (t))
 		   && LAMBDA_TYPE_P (TREE_TYPE (t)))
@@ -14343,6 +14347,14 @@ tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
 	DECL_NAME (r) = omp_reduction_id (ERROR_MARK, DECL_NAME (t),
 					  argtype);
     }
+  else if (DECL_OMP_DECLARE_MAPPER_P (t))
+    {
+      tree argtype
+	= TREE_TYPE (TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (t))));
+      argtype = tsubst (argtype, args, complain, in_decl);
+      if (strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL)
+	DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), argtype);
+    }
 
   if (member && DECL_CONV_FN_P (r))
     /* Type-conversion operator.  Reconstruct the name, in
@@ -18138,6 +18150,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 = omp_instantiate_mappers (new_clauses);
       new_clauses = finish_omp_clauses (new_clauses, ort);
       if (linear_no_step)
 	for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -18877,6 +18891,10 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 			&& DECL_OMP_DECLARE_REDUCTION_P (decl)
 			&& cp_check_omp_declare_reduction (decl))
 		      instantiate_body (pattern_decl, args, decl, true);
+		    else if (TREE_CODE (decl) == FUNCTION_DECL
+			     && DECL_OMP_DECLARE_MAPPER_P (decl)
+			     && cp_check_omp_declare_mapper (decl))
+		      instantiate_body (pattern_decl, args, decl, true);
 		  }
 		else
 		  {
@@ -19827,6 +19845,66 @@ tsubst_omp_udr (tree t, tree args, tsubst_flags_t complain, tree in_decl)
     }
 }
 
+static void
+tsubst_omp_udm (tree t, tree args, tsubst_flags_t complain, tree in_decl)
+{
+  if (t == NULL_TREE || t == error_mark_node)
+    return;
+
+  gcc_assert ((TREE_CODE (t) == STATEMENT_LIST
+	       || TREE_CODE (t) == OMP_DECLARE_MAPPER)
+	      && current_function_decl);
+
+  tree decl = NULL_TREE, mapper;
+
+  /* The function body is:
+
+     statement-list:
+       TYPE t;
+       #pragma omp declare mapper (TYPE t) map(...)
+  */
+
+  if (TREE_CODE (t) == STATEMENT_LIST)
+    {
+      tree_stmt_iterator tsi = tsi_start (t);
+      decl = tsi_stmt (tsi);
+      tsi_next (&tsi);
+      mapper = tsi_stmt (tsi);
+      tsi_next (&tsi);
+      gcc_assert (tsi_end_p (tsi));
+
+      gcc_assert (TREE_CODE (decl) == DECL_EXPR);
+
+      decl = tsubst (DECL_EXPR_DECL (decl), args, complain, in_decl);
+      pushdecl (decl);
+    }
+  else
+    fatal_error (input_location, "malformed OpenMP user-defined mapper");
+
+  if (TREE_CODE (mapper) == DECL_EXPR)
+    mapper = DECL_EXPR_DECL (mapper);
+
+  gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+  tree id = OMP_DECLARE_MAPPER_ID (mapper);
+  tree type = OMP_DECLARE_MAPPER_TYPE (mapper);
+  tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+
+  type = tsubst (type, args, complain, in_decl);
+  /* The _DECLARE_SIMD variant prevents calling finish_omp_clauses on the
+     substituted OMP clauses just yet.  */
+  clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args,
+				complain, in_decl);
+
+  mapper = make_node (OMP_DECLARE_MAPPER);
+  OMP_DECLARE_MAPPER_ID (mapper) = id;
+  OMP_DECLARE_MAPPER_TYPE (mapper) = type;
+  OMP_DECLARE_MAPPER_DECL (mapper) = decl;
+  OMP_DECLARE_MAPPER_CLAUSES (mapper) = clauses;
+  SET_EXPR_LOCATION (mapper, EXPR_LOCATION (t));
+  add_stmt (mapper);
+}
+
 /* T is a postfix-expression that is not being used in a function
    call.  Return the substituted version of T.  */
 
@@ -26602,7 +26680,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
     }
   else
     /* Only OMP reductions are nested.  */
-    gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern));
+    gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern)
+			 || DECL_OMP_DECLARE_MAPPER_P (code_pattern));
 
   vec<tree> omp_privatization_save;
   if (current_function_decl)
@@ -26701,6 +26780,9 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
       if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern))
 	tsubst_omp_udr (DECL_SAVED_TREE (code_pattern), args,
 			tf_warning_or_error, d);
+      else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern))
+	tsubst_omp_udm (DECL_SAVED_TREE (code_pattern), args,
+			tf_warning_or_error, d);
       else
 	{
 	  tsubst_expr (DECL_SAVED_TREE (code_pattern), args,
@@ -26728,6 +26810,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
 
       if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern))
 	cp_check_omp_declare_reduction (d);
+      else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern))
+	cp_check_omp_declare_mapper (d);
     }
 
   /* We're not deferring instantiation any more.  */
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index aa8239c4090..001e2c8eaf9 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -45,6 +45,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "gomp-constants.h"
 #include "predict.h"
 #include "memmodel.h"
+#include "gimplify.h"
 
 /* There routines provide a modular interface to perform many parsing
    operations.  They may therefore be used during actual parsing, or
@@ -4762,7 +4763,8 @@ expand_or_defer_fn_1 (tree fn)
 	   be handled.  */;
       else if (!at_eof
 	       || DECL_IMMEDIATE_FUNCTION_P (fn)
-	       || DECL_OMP_DECLARE_REDUCTION_P (fn))
+	       || DECL_OMP_DECLARE_REDUCTION_P (fn)
+	       || DECL_OMP_DECLARE_MAPPER_P (fn))
 	tentative_decl_linkage (fn);
       else
 	import_export_decl (fn);
@@ -4775,6 +4777,7 @@ expand_or_defer_fn_1 (tree fn)
 	  && !DECL_REALLY_EXTERN (fn)
 	  && !DECL_IMMEDIATE_FUNCTION_P (fn)
 	  && !DECL_OMP_DECLARE_REDUCTION_P (fn)
+	  && !DECL_OMP_DECLARE_MAPPER_P (fn)
 	  && (flag_keep_inline_functions
 	      || (flag_keep_inline_dllexport
 		  && lookup_attribute ("dllexport", DECL_ATTRIBUTES (fn)))))
@@ -4808,7 +4811,8 @@ expand_or_defer_fn_1 (tree fn)
       return false;
     }
 
-  if (DECL_OMP_DECLARE_REDUCTION_P (fn))
+  if (DECL_OMP_DECLARE_REDUCTION_P (fn)
+      || DECL_OMP_DECLARE_MAPPER_P (fn))
     return false;
 
   return true;
@@ -5926,6 +5930,76 @@ omp_reduction_lookup (location_t loc, tree id, tree type, tree *baselinkp,
   return id;
 }
 
+/* Return identifier to look up for omp declare mapper.  */
+
+tree
+omp_mapper_id (tree mapper_id, tree type)
+{
+  const char *p = NULL;
+  const char *m = NULL;
+
+  if (mapper_id == NULL_TREE)
+    p = "";
+  else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE)
+    p = IDENTIFIER_POINTER (mapper_id);
+  else
+    return error_mark_node;
+
+  if (type != NULL_TREE)
+    m = mangle_type_string (TYPE_MAIN_VARIANT (type));
+
+  const char prefix[] = "omp declare mapper ";
+  size_t lenp = sizeof (prefix);
+  if (strncmp (p, prefix, lenp - 1) == 0)
+    lenp = 1;
+  size_t len = strlen (p);
+  size_t lenm = m ? strlen (m) + 1 : 0;
+  char *name = XALLOCAVEC (char, lenp + len + lenm);
+  memcpy (name, prefix, lenp - 1);
+  memcpy (name + lenp - 1, p, len + 1);
+  if (m)
+    {
+      name[lenp + len - 1] = '~';
+      memcpy (name + lenp + len, m, lenm);
+    }
+  return get_identifier (name);
+}
+
+static tree
+omp_mapper_lookup (tree id, tree type)
+{
+  if (TREE_CODE (type) != RECORD_TYPE
+      && TREE_CODE (type) != UNION_TYPE)
+    return NULL_TREE;
+  id = omp_mapper_id (id, type);
+  return lookup_name (id);
+}
+
+static tree
+omp_extract_mapper_directive (tree fndecl)
+{
+  if (BASELINK_P (fndecl))
+    /* See through BASELINK nodes to the underlying function.  */
+    fndecl = BASELINK_FUNCTIONS (fndecl);
+
+  tree body = DECL_SAVED_TREE (fndecl);
+
+  if (TREE_CODE (body) == BIND_EXPR)
+    body = BIND_EXPR_BODY (body);
+
+  if (TREE_CODE (body) == STATEMENT_LIST)
+    {
+      tree_stmt_iterator tsi = tsi_start (body);
+      gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR);
+      tsi_next (&tsi);
+      body = tsi_stmt (tsi);
+    }
+
+  gcc_assert (TREE_CODE (body) == OMP_DECLARE_MAPPER);
+
+  return body;
+}
+
 /* Helper function for cp_parser_omp_declare_reduction_exprs
    and tsubst_omp_udr.
    Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6407,6 +6481,31 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor)
   return false;
 }
 
+/* Check an instance of an "omp declare mapper" function.  */
+
+bool
+cp_check_omp_declare_mapper (tree udm)
+{
+  tree type = TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (udm)));
+  location_t loc = DECL_SOURCE_LOCATION (udm);
+  gcc_assert (TYPE_REF_P (type));
+  type = TREE_TYPE (type);
+
+  if (type == error_mark_node)
+    return false;
+
+  if (!processing_template_decl
+      && TREE_CODE (type) != RECORD_TYPE
+      && TREE_CODE (type) != UNION_TYPE)
+    {
+      error_at (loc, "%qT is not a struct, union or class type in "
+		"%<#pragma omp declare mapper%>", type);
+      return false;
+    }
+
+  return true;
+}
+
 /* Called from finish_struct_1.  linear(this) or linear(this:step)
    clauses might not be finalized yet because the class has been incomplete
    when parsing #pragma omp declare simd methods.  Fix those up now.  */
@@ -6667,6 +6766,242 @@ cp_oacc_check_attachments (tree c)
   return false;
 }
 
+struct remap_mapper_decl_info
+{
+  tree dummy_var;
+  tree expr;
+};
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+  remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data;
+
+  if (operand_equal_p (*tp, map_info->dummy_var))
+    {
+      *tp = map_info->expr;
+      *walk_subtrees = 0;
+    }
+
+  return NULL_TREE;
+}
+
+static tree *
+omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+			enum gomp_map_kind outer_kind)
+{
+  tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+  tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+  tree mapper_name = NULL_TREE;
+
+  remap_mapper_decl_info map_info;
+  map_info.dummy_var = dummy_var;
+  map_info.expr = expr;
+
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      tree unshared = unshare_expr (c);
+      enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c);
+      tree t = OMP_CLAUSE_DECL (unshared);
+      tree type = NULL_TREE;
+      bool nonunit_array_with_mapper = false;
+
+      if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+	{
+	  mapper_name = t;
+	  continue;
+	}
+      else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+	{
+	  mapper_name = NULL_TREE;
+	  continue;
+	}
+
+      if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+	{
+	  tree low = TREE_OPERAND (t, 1);
+	  tree len = TREE_OPERAND (t, 2);
+
+	  if (len && integer_onep (len))
+	    {
+	      t = TREE_OPERAND (t, 0);
+
+	      if (POINTER_TYPE_P (TREE_TYPE (t))
+		  || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+		type = TREE_TYPE (TREE_TYPE (t));
+
+	      if (!low)
+		low = integer_zero_node;
+
+	      if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
+		t = convert_from_reference (t);
+
+	      t = build_array_ref (OMP_CLAUSE_LOCATION (c), t, low);
+	    }
+	  else
+	    {
+	      type = TREE_TYPE (t);
+	      nonunit_array_with_mapper = true;
+	    }
+	}
+      else
+	type = TREE_TYPE (t);
+
+      gcc_assert (type);
+
+      if (type == error_mark_node)
+	continue;
+
+      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);
+
+      type = TYPE_MAIN_VARIANT (type);
+
+      tree mapper_fn = omp_mapper_lookup (mapper_name, type);
+
+      if (mapper_fn && nonunit_array_with_mapper)
+	{
+	  sorry ("user-defined mapper with non-unit length array section");
+	  continue;
+	}
+      else if (mapper_fn)
+	{
+	  tree nested_mapper = 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);
+	      continue;
+	    }
+	}
+      else if (mapper_name)
+	{
+	  error ("mapper %qE not found for type %qT", mapper_name, type);
+	  continue;
+	}
+
+      *outlist = unshared;
+      outlist = &OMP_CLAUSE_CHAIN (unshared);
+    }
+
+  return outlist;
+}
+
+tree
+omp_instantiate_mappers (tree clauses)
+{
+  tree c, *pc, mapper_name = NULL_TREE;
+
+  for (pc = &clauses, c = clauses; c; c = *pc)
+    {
+      bool using_mapper = false;
+
+      switch (OMP_CLAUSE_CODE (c))
+	{
+	case OMP_CLAUSE_MAP:
+	  {
+	    tree t = OMP_CLAUSE_DECL (c);
+	    tree type = NULL_TREE;
+	    bool nonunit_array_with_mapper = false;
+
+	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+	      {
+		if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
+		  mapper_name = OMP_CLAUSE_DECL (c);
+		else
+		  mapper_name = NULL_TREE;
+		pc = &OMP_CLAUSE_CHAIN (c);
+		continue;
+	      }
+
+	    gcc_assert (TREE_CODE (t) != TREE_LIST);
+
+	    if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+	      {
+		tree low = TREE_OPERAND (t, 1);
+		tree len = TREE_OPERAND (t, 2);
+
+		if (len && integer_onep (len))
+		  {
+		    t = TREE_OPERAND (t, 0);
+
+		    if (!TREE_TYPE (t))
+		      {
+			pc = &OMP_CLAUSE_CHAIN (c);
+			continue;
+		      }
+
+		    if (POINTER_TYPE_P (TREE_TYPE (t))
+			|| TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+		      type = TREE_TYPE (TREE_TYPE (t));
+
+		    if (!low)
+		      low = integer_zero_node;
+		  }
+		else
+		  {
+		    /* !!! Array sections of size >1 with mappers for elements
+		       are hard to support.  Do something here.  */
+		    nonunit_array_with_mapper = true;
+		    type = TREE_TYPE (t);
+		  }
+	      }
+	    else
+	      type = TREE_TYPE (t);
+
+	    if (type == NULL_TREE || type == error_mark_node)
+	      {
+		pc = &OMP_CLAUSE_CHAIN (c);
+		continue;
+	      }
+
+	    enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
+	    if (kind == GOMP_MAP_UNSET)
+	      kind = GOMP_MAP_TOFROM;
+
+	    type = TYPE_MAIN_VARIANT (type);
+
+	    tree mapper_fn = omp_mapper_lookup (mapper_name, type);
+
+	    if (mapper_fn && nonunit_array_with_mapper)
+	      {
+		sorry ("user-defined mapper with non-unit length "
+		       "array section");
+		using_mapper = true;
+	      }
+	    else if (mapper_fn)
+	      {
+		tree mapper = omp_extract_mapper_directive (mapper_fn);
+		pc = omp_instantiate_mapper (pc, mapper, t, kind);
+		using_mapper = true;
+	      }
+	    else if (mapper_name)
+	      {
+		error ("mapper %qE not found for type %qT", mapper_name, type);
+		using_mapper = true;
+	      }
+	  }
+	  break;
+
+	default:
+	  ;
+	}
+
+      if (using_mapper)
+	*pc = OMP_CLAUSE_CHAIN (c);
+      else
+	pc = &OMP_CLAUSE_CHAIN (c);
+    }
+
+  return clauses;
+}
+
 /* For all elements of CLAUSES, validate them vs OpenMP constraints.
    Remove any elements from the list that are invalid.  */
 
@@ -7949,6 +8284,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_MAP:
 	  if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
 	    goto move_implicit;
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+	    {
+	      remove = true;
+	      break;
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -9378,6 +9719,108 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
   return add_stmt (stmt);
 }
 
+struct mapper_list
+{
+  hash_set<omp_name_type> *seen_types;
+  vec<tree> *mappers;
+
+  mapper_list (hash_set<omp_name_type> *s, vec<tree> *m)
+    : seen_types (s), mappers (m) { }
+
+  void add_mapper (tree name, tree type, tree mapperfn)
+  {
+    /* We can't hash a NULL_TREE...  */
+    if (!name)
+      name = void_node;
+
+    omp_name_type n_t = { name, type };
+
+    if (seen_types->contains (n_t))
+      return;
+
+    seen_types->add (n_t);
+    mappers->safe_push (mapperfn);
+  }
+
+  bool contains (tree name, tree type)
+  {
+    if (!name)
+      name = void_node;
+
+    return seen_types->contains ({ name, type });
+  }
+};
+
+static void
+find_nested_mappers (mapper_list *mlist, tree mapper_fn)
+{
+  tree mapper = omp_extract_mapper_directive (mapper_fn);
+  tree mapper_name = NULL_TREE;
+
+  if (mapper == error_mark_node)
+    return;
+
+  gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+  for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+       clause;
+       clause = OMP_CLAUSE_CHAIN (clause))
+    {
+      tree expr = OMP_CLAUSE_DECL (clause);
+      enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
+      tree elem_type;
+
+      if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+	{
+	  mapper_name = expr;
+	  continue;
+	}
+      else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+	{
+	  mapper_name = NULL_TREE;
+	  continue;
+	}
+
+      gcc_assert (TREE_CODE (expr) != TREE_LIST);
+      if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+	{
+	  while (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+	    expr = TREE_OPERAND (expr, 0); //TREE_CHAIN (expr);
+
+	  elem_type = TREE_TYPE (expr);
+	}
+      else
+	elem_type = TREE_TYPE (expr);
+
+      /* This might be too much... or not enough?  */
+      while (TREE_CODE (elem_type) == ARRAY_TYPE
+	     || TREE_CODE (elem_type) == POINTER_TYPE
+	     || TREE_CODE (elem_type) == REFERENCE_TYPE)
+	elem_type = TREE_TYPE (elem_type);
+
+      elem_type = TYPE_MAIN_VARIANT (elem_type);
+
+      if (AGGREGATE_TYPE_P (elem_type)
+	  && !mlist->contains (mapper_name, elem_type))
+	{
+	  tree nested_mapper_fn
+	    = omp_mapper_lookup (mapper_name, elem_type);
+
+	  if (nested_mapper_fn)
+	    {
+	      mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
+	      find_nested_mappers (mlist, nested_mapper_fn);
+	    }
+	  else if (mapper_name)
+	    {
+	      error ("mapper %qE not found for type %qT", mapper_name,
+		     elem_type);
+	      continue;
+	    }
+	}
+    }
+}
+
 /* Used to walk OpenMP target directive body.  */
 
 struct omp_target_walk_data
@@ -9403,6 +9846,8 @@ struct omp_target_walk_data
   /* Local variables declared inside a BIND_EXPR, used to filter out such
      variables when recording lambda_objects_accessed.  */
   hash_set<tree> local_decls;
+
+  mapper_list *mappers;
 };
 
 /* Helper function of finish_omp_target_clauses, called via
@@ -9416,6 +9861,8 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
   struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;
   tree current_object = data->current_object;
   tree current_closure = data->current_closure;
+  mapper_list *mlist = data->mappers;
+  tree aggr_type = NULL_TREE;
 
   /* References inside of these expression codes shouldn't incur any
      form of mapping, so return early.  */
@@ -9429,6 +9876,22 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
   if (TREE_CODE (t) == OMP_CLAUSE)
     return NULL_TREE;
 
+  if (TREE_CODE (t) == COMPONENT_REF
+      && AGGREGATE_TYPE_P (TREE_TYPE (TREE_OPERAND (t, 0))))
+    aggr_type = TREE_TYPE (TREE_OPERAND (t, 0));
+  else if ((TREE_CODE (t) == VAR_DECL
+	    || TREE_CODE (t) == PARM_DECL
+	    || TREE_CODE (t) == RESULT_DECL)
+	   && AGGREGATE_TYPE_P (TREE_TYPE (t)))
+    aggr_type = TREE_TYPE (t);
+
+  if (aggr_type)
+    {
+      tree mapper_fn = omp_mapper_lookup (NULL_TREE, aggr_type);
+      if (mapper_fn)
+	mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn);
+    }
+
   if (current_object)
     {
       tree this_expr = TREE_OPERAND (current_object, 0);
@@ -9532,10 +9995,38 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
   else
     data.current_closure = NULL_TREE;
 
+  hash_set<omp_name_type> seen_types;
+  auto_vec<tree> mapper_fns;
+  mapper_list mlist (&seen_types, &mapper_fns);
+  data.mappers = &mlist;
+
   cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
 
+  unsigned int i;
+  tree mapper_fn;
+  FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
+    find_nested_mappers (&mlist, mapper_fn);
+
   auto_vec<tree, 16> new_clauses;
 
+  FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
+    {
+      tree mapper = omp_extract_mapper_directive (mapper_fn);
+      if (mapper == error_mark_node)
+	continue;
+      tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper);
+      tree decl = OMP_DECLARE_MAPPER_DECL (mapper);
+      if (BASELINK_P (mapper_fn))
+	mapper_fn = BASELINK_FUNCTIONS (mapper_fn);
+
+      tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_);
+      OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name;
+      OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl;
+      OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper_fn;
+
+      new_clauses.safe_push (c);
+    }
+
   tree omp_target_this_expr = NULL_TREE;
   tree *explicit_this_deref_map = NULL;
   if (data.this_expr_accessed)
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 80492c952aa..34245e106b7 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -27,6 +27,9 @@ along with GCC; see the file COPYING3.  If not see
 #include "match.h"
 #include "parse.h"
 #include "tree-core.h"
+#include "tree.h"
+#include "fold-const.h"
+#include "tree-hash-traits.h"
 #include "omp-general.h"
 
 /* Current statement label.  Zero means no statement label.  Because new_st
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index a08e3468788..e14efe780cf 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -219,6 +219,7 @@ struct gimplify_omp_ctx
 {
   struct gimplify_omp_ctx *outer_context;
   splay_tree variables;
+  hash_map<omp_name_type, tree> *implicit_mappers;
   hash_set<tree> *privatized_types;
   tree clauses;
   /* Iteration variables in an OMP_FOR.  */
@@ -452,6 +453,7 @@ new_omp_context (enum omp_region_type region_type)
   c = XCNEW (struct gimplify_omp_ctx);
   c->outer_context = gimplify_omp_ctxp;
   c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
+  c->implicit_mappers = new hash_map<omp_name_type, tree>;
   c->privatized_types = new hash_set<tree>;
   c->location = input_location;
   c->region_type = region_type;
@@ -475,6 +477,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
 {
   splay_tree_delete (c->variables);
   delete c->privatized_types;
+  delete c->implicit_mappers;
   c->loop_iter_var.release ();
   XDELETE (c);
 }
@@ -11328,21 +11331,190 @@ error_out:
   return success;
 }
 
-static void
-gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
-			   enum omp_region_type region_type,
-			   enum tree_code code)
+struct instantiate_mapper_info
 {
-  using namespace omp_addr_tokenizer;
-  struct gimplify_omp_ctx *ctx, *outer_ctx;
-  tree c;
-  tree *prev_list_p = NULL, *orig_list_p = list_p;
-  int handled_depend_iterators = -1;
-  int nowait = -1;
+  hash_set<tree> *handled_structs;
+  tree *mapper_clauses_p;
+  struct gimplify_omp_ctx *omp_ctx;
+};
 
-  ctx = new_omp_context (region_type);
+struct remap_mapper_decl_info
+{
+  tree dummy_var;
+  tree expr;
+};
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+  remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data;
+
+  if (operand_equal_p (*tp, map_info->dummy_var))
+    {
+      *tp = unshare_expr (map_info->expr);
+      *walk_subtrees = 0;
+    }
+
+  return NULL_TREE;
+}
+
+static tree *
+omp_instantiate_mapper (hash_map<omp_name_type, tree> *implicit_mappers,
+			tree mapper, tree expr, enum gomp_map_kind outer_kind,
+			tree *mapper_clauses_p)
+{
+  tree mapper_name = NULL_TREE;
+  gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+  tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+  tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+
+  remap_mapper_decl_info map_info;
+
+  map_info.dummy_var = dummy_var;
+  map_info.expr = expr;
+
+  for (; clause; clause = OMP_CLAUSE_CHAIN (clause))
+    {
+      enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
+
+      if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+	{
+	  mapper_name = OMP_CLAUSE_DECL (clause);
+	  continue;
+	}
+      else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+	{
+	  mapper_name = NULL_TREE;
+	  continue;
+	}
+
+      tree decl = OMP_CLAUSE_DECL (clause), unshared;
+
+      if (TREE_CODE (decl) == OMP_ARRAY_SECTION
+	  && TREE_OPERAND (decl, 2)
+	  && integer_onep (TREE_OPERAND (decl, 2)))
+	{
+	  unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+				       OMP_CLAUSE_CODE (clause));
+	  tree low = TREE_OPERAND (decl, 1);
+	  if (!low || integer_zerop (low))
+	    OMP_CLAUSE_DECL (unshared)
+	      = build_fold_indirect_ref (TREE_OPERAND (decl, 0));
+	  else
+	    OMP_CLAUSE_DECL (unshared) = decl;
+	  OMP_CLAUSE_SIZE (unshared) = OMP_CLAUSE_SIZE (clause);
+	}
+      else
+	unshared = unshare_expr (clause);
+
+      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);
+
+      decl = OMP_CLAUSE_DECL (unshared);
+      tree type = TYPE_MAIN_VARIANT (TREE_TYPE (decl));
+
+      tree *nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
+
+      if (nested_mapper_p && *nested_mapper_p != mapper)
+	{
+	  if (clause_kind == GOMP_MAP_UNSET)
+	    clause_kind = outer_kind;
+
+	  mapper_clauses_p
+	    = omp_instantiate_mapper (implicit_mappers, *nested_mapper_p,
+				      decl, clause_kind, mapper_clauses_p);
+	  continue;
+	}
+
+      *mapper_clauses_p = unshared;
+      mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared);
+    }
+
+  return mapper_clauses_p;
+}
+
+/* Scan GROUPS for mappings that map all or part of a struct or union decl,
+   and add to HANDLED.  Implicit user-defined mappers will then not be invoked
+   for those decls.  */
+
+static void
+omp_find_explicitly_mapped_structs (hash_set<tree> *handled,
+				    vec<omp_mapping_group> *groups)
+{
+  if (!groups)
+    return;
+
+  for (auto &i : *groups)
+    {
+      tree node = *i.grp_start;
+
+      if (OMP_CLAUSE_CODE (node) == OMP_CLAUSE_MAP
+	  && DECL_P (OMP_CLAUSE_DECL (node)))
+	switch (OMP_CLAUSE_MAP_KIND (node) & ~GOMP_MAP_FLAG_FORCE)
+	  {
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_STRUCT:
+	    handled->add (OMP_CLAUSE_DECL (node));
+	    break;
+	  default:
+	    ;
+	  }
+    }
+}
+
+static int
+omp_instantiate_implicit_mappers (splay_tree_node n, void *data)
+{
+  tree decl = (tree) n->key;
+  instantiate_mapper_info *im_info = (instantiate_mapper_info *) data;
+  gimplify_omp_ctx *ctx = im_info->omp_ctx;
+  tree *mapper_p = NULL;
+  tree type = TREE_TYPE (decl);
+  bool ref_p = false;
+
+  if (TREE_CODE (type) == REFERENCE_TYPE)
+    {
+      ref_p = true;
+      type = TREE_TYPE (type);
+    }
+
+  type = TYPE_MAIN_VARIANT (type);
+
+  if (DECL_P (decl) && type && AGGREGATE_TYPE_P (type))
+    {
+      gcc_assert (ctx);
+      mapper_p = ctx->implicit_mappers->get ({ NULL_TREE, type });
+    }
+
+  bool handled_p = im_info->handled_structs->contains (decl);
+
+  if (mapper_p && !handled_p)
+    {
+      /* If we have a reference, map the pointed-to object rather than the
+	 reference itself.  */
+      if (ref_p)
+	decl = build_fold_indirect_ref (decl);
+
+      im_info->mapper_clauses_p
+	= omp_instantiate_mapper (ctx->implicit_mappers, *mapper_p, decl,
+				  GOMP_MAP_TOFROM, im_info->mapper_clauses_p);
+    }
+
+  return 0;
+}
+
+static struct gimplify_omp_ctx *
+new_omp_context_for_scan (enum omp_region_type region_type,
+			  enum tree_code code)
+{
+  struct gimplify_omp_ctx *ctx = new_omp_context (region_type);
   ctx->code = code;
-  outer_ctx = ctx->outer_context;
   if (code == OMP_TARGET)
     {
       if (!lang_GNU_Fortran ())
@@ -11366,19 +11538,48 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       default:
 	break;
       }
+  gimplify_omp_ctxp = ctx;
+  return ctx;
+}
 
-  if (code == OMP_TARGET
-      || code == OMP_TARGET_DATA
-      || code == OMP_TARGET_ENTER_DATA
-      || code == OMP_TARGET_EXIT_DATA)
+/* Scan the OMP clauses in *LIST_P, installing mappings into a new
+   and previous omp contexts.  */
+
+static void
+gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
+			   enum omp_region_type region_type,
+			   enum tree_code code,
+			   struct gimplify_omp_ctx *existing_ctx = NULL,
+			   bool suppress_gimplification = false)
+{
+  struct gimplify_omp_ctx *ctx, *outer_ctx;
+  tree c;
+  tree *prev_list_p = NULL, *orig_list_p = list_p;
+  int handled_depend_iterators = -1;
+  int nowait = -1;
+  bool second_pass = existing_ctx != NULL && !suppress_gimplification;
+
+  if (existing_ctx)
+    ctx = existing_ctx;
+  else
+    ctx = new_omp_context_for_scan (region_type, code);
+
+  outer_ctx = ctx->outer_context;
+  gimplify_omp_ctxp = outer_ctx;
+
+  if (!suppress_gimplification
+      && (code == OMP_TARGET
+	  || code == OMP_TARGET_DATA
+	  || code == OMP_TARGET_ENTER_DATA
+	  || code == OMP_TARGET_EXIT_DATA))
     {
       vec<omp_mapping_group> *groups;
       groups = omp_gather_mapping_groups (list_p);
+      hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+
       if (groups)
 	{
-	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
 	  grpmap = omp_index_mapping_groups (groups);
-
 	  omp_resolve_clause_dependencies (code, groups, grpmap);
 	  omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
 					  list_p);
@@ -11433,6 +11634,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       tree decl;
       auto_vec<omp_addr_token *, 10> addr_tokens;
 
+      if (second_pass && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+	{
+	  list_p = &OMP_CLAUSE_CHAIN (c);
+	  continue;
+	}
+
       switch (OMP_CLAUSE_CODE (c))
 	{
 	case OMP_CLAUSE_PRIVATE:
@@ -11834,8 +12041,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
-			     NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+	  if (!suppress_gimplification
+	      && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
+				NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
 	      remove = true;
 	      break;
@@ -11854,9 +12062,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
 
-	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-	      && (addr_tokens[0]->type == STRUCTURE_BASE
-		  || addr_tokens[0]->type == ARRAY_BASE)
+	  /* Adding the decl for a struct access: if we're suppressing
+	     gimplification (on the first pass before "declare mapper"
+	     instantiation), we haven't built sibling lists yet, so process
+	     decls that will be turned into GOMP_MAP_STRUCT in the second
+	     pass.  */
+	  if (((!suppress_gimplification
+		&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+		&& (addr_tokens[0]->type == STRUCTURE_BASE
+		    || addr_tokens[0]->type == ARRAY_BASE))
+	       || (suppress_gimplification
+		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+		   && addr_tokens[0]->type == STRUCTURE_BASE))
 	      && addr_tokens[0]->u.structure_base_kind == BASE_DECL)
 	    {
 	      gcc_assert (addr_tokens[1]->type == ACCESS_METHOD);
@@ -11865,6 +12082,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      if (omp_access_chain_p (addr_tokens, 1))
 		break;
 	      decl = addr_tokens[1]->expr;
+	      if (suppress_gimplification
+		  && splay_tree_lookup (ctx->variables, (splay_tree_key) decl))
+		break;
 	      flags = GOVD_MAP | GOVD_EXPLICIT;
 
 	      gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT
@@ -11874,9 +12094,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	  if (TREE_CODE (decl) == TARGET_EXPR)
 	    {
-	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
-				 is_gimple_lvalue, fb_lvalue)
-		  == GS_ERROR)
+	      if (!suppress_gimplification
+		  && (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+				     is_gimple_lvalue, fb_lvalue)
+		      == GS_ERROR))
 		remove = true;
 	    }
 	  else if (!DECL_P (decl))
@@ -11954,7 +12175,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		    }
 		}
 
-	      if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
+	      if (!suppress_gimplification
+		  && code == OMP_TARGET
+		  && OMP_CLAUSE_MAP_IN_REDUCTION (c))
 		{
 		  /* Don't gimplify *pd fully at this point, as the base
 		     will need to be adjusted during omp lowering.  */
@@ -12062,8 +12285,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		    }
 		}
-	      else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
-				      fb_lvalue) == GS_ERROR)
+	      else if (!suppress_gimplification
+		       && gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+					 fb_lvalue) == GS_ERROR)
 		{
 		  remove = true;
 		  break;
@@ -12213,6 +12437,29 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  goto do_notice;
 
+	case OMP_CLAUSE__MAPPER_BINDING_:
+	  {
+	    tree name = OMP_CLAUSE__MAPPER_BINDING__ID (c);
+	    tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c);
+	    tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var));
+	    tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c);
+	    tree mapper = DECL_SAVED_TREE (fndecl);
+	    if (TREE_CODE (mapper) == BIND_EXPR)
+	      mapper = BIND_EXPR_BODY (mapper);
+	    if (TREE_CODE (mapper) == STATEMENT_LIST)
+	      {
+		tree_stmt_iterator tsi = tsi_start (mapper);
+		gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR);
+		tsi_next (&tsi);
+		mapper = tsi_stmt (tsi);
+	      }
+	    gcc_assert (mapper != NULL_TREE
+			&& TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+	    ctx->implicit_mappers->put ({ name, type }, mapper);
+	    remove = true;
+	    break;
+	  }
+
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  flags = GOVD_EXPLICIT;
@@ -12760,7 +13007,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
-  if ((region_type & ORT_TARGET) != 0)
+  if ((region_type & ORT_TARGET) != 0 && !suppress_gimplification)
     /* If we have a target region, we can push all the attaches to the end of
        the list (we may have standalone "attach" operations synthesized for
        GOMP_MAP_STRUCT nodes that must be processed after the attachment point
@@ -15984,10 +16231,25 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     }
 
   bool save_in_omp_construct = in_omp_construct;
+
   if ((ort & ORT_ACC) == 0)
     in_omp_construct = false;
-  gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
-			     TREE_CODE (expr));
+
+  struct gimplify_omp_ctx *ctx = NULL;
+
+  if (TREE_CODE (expr) == OMP_TARGET
+      || TREE_CODE (expr) == OMP_TARGET_DATA
+      || TREE_CODE (expr) == OMP_TARGET_ENTER_DATA
+      || TREE_CODE (expr) == OMP_TARGET_EXIT_DATA)
+    {
+      ctx = new_omp_context_for_scan (ort, TREE_CODE (expr));
+      gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+				 TREE_CODE (expr), ctx, true);
+    }
+  else
+    gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+			       TREE_CODE (expr));
+
   if (TREE_CODE (expr) == OMP_TARGET)
     optimize_target_teams (expr, pre_p);
   if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -16025,6 +16287,45 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     }
   else
     gimplify_and_add (OMP_BODY (expr), &body);
+  if (TREE_CODE (expr) == OMP_TARGET
+      || TREE_CODE (expr) == OMP_TARGET_DATA
+      || TREE_CODE (expr) == OMP_TARGET_ENTER_DATA
+      || TREE_CODE (expr) == OMP_TARGET_EXIT_DATA)
+    {
+      tree mapper_clauses = NULL_TREE;
+      instantiate_mapper_info im_info;
+      hash_set<tree> handled_structs;
+
+      /* If we already have clauses pertaining to a struct variable, then
+	 we don't want to implicitly invoke a user-defined mapper.  Scan
+	 through the groups to check what we have already.  */
+
+      splay_tree_foreach (ctx->variables,
+			  omp_find_explicitly_mapped_structs_r,
+			  (void *) &handled_structs);
+
+      im_info.handled_structs = &handled_structs;
+      im_info.mapper_clauses_p = &mapper_clauses;
+      im_info.omp_ctx = ctx;
+      im_info.pre_p = pre_p;
+
+      splay_tree_foreach (ctx->variables,
+			  omp_instantiate_implicit_mappers,
+			  (void *) &im_info);
+
+      if (mapper_clauses)
+	mapper_clauses
+	  = lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses);
+
+      /* Stick the implicitly-expanded mapper clauses at the end of the
+	 clause list.  */
+      tree *tail = &OMP_CLAUSES (expr);
+      while (*tail)
+	tail = &OMP_CLAUSE_CHAIN (*tail);
+      *tail = mapper_clauses;
+      gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+				 TREE_CODE (expr), ctx);
+    }
   gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr),
 			       TREE_CODE (expr));
   in_omp_construct = save_in_omp_construct;
@@ -16684,6 +16985,16 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
   return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
 }
 
+static enum gimplify_status
+gimplify_omp_declare_mapper (tree *expr_p)
+{
+  /* We don't want assembler output -- this inhibits it.  */
+  DECL_DECLARED_INLINE_P (current_function_decl) = 1;
+
+  *expr_p = NULL_TREE;
+  return GS_ALL_DONE;
+}
+
 /* Convert the GENERIC expression tree *EXPR_P to GIMPLE.  If the
    expression produces a value to be used as an operand inside a GIMPLE
    statement, the value will be stored back in *EXPR_P.  This value will
@@ -17603,6 +17914,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = gimplify_omp_atomic (expr_p, pre_p);
 	  break;
 
+	case OMP_DECLARE_MAPPER:
+	  ret = gimplify_omp_declare_mapper (expr_p);
+	  break;
+
 	case TRANSACTION_EXPR:
 	  ret = gimplify_transaction (expr_p, pre_p);
 	  break;
diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h
index 0f1bfbf255c..dcb1e8c383e 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -85,6 +85,7 @@ extern enum omp_clause_defaultmap_kind lhd_omp_predetermined_mapping (tree);
 extern tree lhd_omp_assignment (tree, tree, tree);
 extern void lhd_omp_finish_clause (tree, gimple_seq *, bool);
 extern tree lhd_omp_array_size (tree, gimple_seq *);
+extern tree lhd_omp_finish_mapper_clauses (tree);
 struct gimplify_omp_ctx;
 extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
@@ -273,6 +274,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
 #define LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR NULL
 #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
 #define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES lhd_omp_finish_mapper_clauses
 #define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false
 #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
 #define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false
@@ -307,6 +309,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
   LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \
   LANG_HOOKS_OMP_CLAUSE_DTOR, \
   LANG_HOOKS_OMP_FINISH_CLAUSE, \
+  LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, \
   LANG_HOOKS_OMP_ALLOCATABLE_P, \
   LANG_HOOKS_OMP_SCALAR_P, \
   LANG_HOOKS_OMP_SCALAR_TARGET_P, \
diff --git a/gcc/langhooks.cc b/gcc/langhooks.cc
index 7d975675990..7960fe05b24 100644
--- a/gcc/langhooks.cc
+++ b/gcc/langhooks.cc
@@ -642,6 +642,15 @@ lhd_omp_array_size (tree, gimple_seq *)
   return NULL_TREE;
 }
 
+/* Finalize clause list C after expanding custom mappers for implicitly-mapped
+   variables.  */
+
+tree
+lhd_omp_finish_mapper_clauses (tree c)
+{
+  return c;
+}
+
 /* Return true if DECL is a scalar variable (for the purpose of
    implicit firstprivatization & mapping). Only if alloc_ptr_ok
    are allocatables and pointers accepted. */
diff --git a/gcc/langhooks.h b/gcc/langhooks.h
index b1b2b0e10f0..8ccbab83db3 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -313,6 +313,10 @@ struct lang_hooks_for_decls
   /* Do language specific checking on an implicitly determined clause.  */
   void (*omp_finish_clause) (tree clause, gimple_seq *pre_p, bool);
 
+  /* Finish language-specific processing on mapping nodes after expanding
+     user-defined mappers.  */
+  tree (*omp_finish_mapper_clauses) (tree clauses);
+
   /* Return true if DECL is an allocatable variable (for the purpose of
      implicit mapping).  */
   bool (*omp_allocatable_p) (tree decl);
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 6fc0b11996a..80325e5e4d7 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -209,4 +209,56 @@ typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
 
 extern bool omp_parse_expr (vec<omp_addr_token *> &, tree);
 
+struct omp_name_type
+{
+  tree name;
+  tree type;
+};
+
+template <>
+struct default_hash_traits <omp_name_type>
+  : typed_noop_remove <omp_name_type>
+{
+  GTY((skip)) typedef omp_name_type value_type;
+  GTY((skip)) typedef omp_name_type compare_type;
+
+  static hashval_t
+  hash (omp_name_type p)
+  {
+    return p.name ? iterative_hash_expr (p.name, TYPE_UID (p.type))
+		  : TYPE_UID (p.type);
+  }
+
+  static const bool empty_zero_p = true;
+
+  static bool
+  is_empty (omp_name_type p)
+  {
+    return p.type == NULL;
+  }
+
+  static bool
+  is_deleted (omp_name_type)
+  {
+    return false;
+  }
+
+  static bool
+  equal (const omp_name_type &a, const omp_name_type &b)
+  {
+    if (a.name == NULL_TREE && b.name == NULL_TREE)
+      return a.type == b.type;
+    else if (a.name == NULL_TREE || b.name == NULL_TREE)
+      return false;
+    else
+      return a.name == b.name && a.type == b.type;
+  }
+
+  static void
+  mark_empty (omp_name_type &e)
+  {
+    e.type = NULL;
+  }
+};
+
 #endif /* GCC_OMP_GENERAL_H */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index c749db845b0..77054b6b56a 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,10 +13,12 @@ foo (void)
   #pragma omp target map (to:a)
   ;
 
-  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
   ;
 
-  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
   ;
 
   #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
new file mode 100644
index 00000000000..3177d20adbc
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
@@ -0,0 +1,58 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+  S s;
+  s.ptr = new int[N];
+  s.size = N;
+
+#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size])
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(default), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(foo), alloc: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(bar), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+  return 0;
+}
+
+// { 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" } }
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C
new file mode 100644
index 00000000000..06d999ea654
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C
@@ -0,0 +1,30 @@
+// { dg-do compile }
+
+// Error-checking tests for "omp declare mapper".
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+struct Z {
+  int z;
+};
+
+int main (int argc, char *argv[])
+{
+#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size])
+
+  /* This one's a duplicate.  */
+#pragma omp declare mapper (default: S v) map (to: v.size) map (v) // { dg-error "redeclaration of 'pragma omp declare mapper'" }
+
+  /* ...and this one doesn't use a "base language identifier" for the mapper
+     name.  */
+#pragma omp declare mapper (case: S v) map (to: v.size) // { dg-error "expected identifier or 'default' before 'case'" }
+  // { dg-error "expected ':' before 'case'" "" { target *-*-* } .-1 }
+
+  /* A non-struct/class/union type isn't supposed to work.  */
+#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z) // { dg-error "'Z \\\[5\\\]' is not a struct, union or class type in '#pragma omp declare mapper'" }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C
new file mode 100644
index 00000000000..92212fd0dbd
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-3.C
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// Test named mapper invocation.
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+int main (int argc, char *argv[])
+{
+  int N = 1024;
+#pragma omp declare mapper (mapN:S s) map(to:s.ptr, s.size) map(s.ptr[:N])
+
+  S s;
+  s.ptr = new int[N];
+
+#pragma omp target map(mapper(mapN), tofrom: s)
+// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(to:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} "gimple" } }
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C
new file mode 100644
index 00000000000..85bef470332
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-4.C
@@ -0,0 +1,74 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+// Check mapper binding clauses.
+
+struct Y {
+  int z;
+};
+
+struct Z {
+  int z;
+};
+
+#pragma omp declare mapper (Y y) map(tofrom: y)
+#pragma omp declare mapper (Z z) map(tofrom: z)
+
+int foo (void)
+{
+  Y yy;
+  Z zz;
+  int dummy;
+
+#pragma omp target data map(dummy)
+  {
+  #pragma omp target
+    {
+      yy.z++;
+      zz.z++;
+    }
+    yy.z++;
+  }
+  return yy.z;
+}
+
+struct P
+{
+  Z *zp;
+};
+
+int bar (void)
+{
+  Y yy;
+  Z zz;
+  P pp;
+  Z t;
+  int dummy;
+
+  pp.zp = &t;
+
+#pragma omp declare mapper (Y y) map(tofrom: y.z)
+#pragma omp declare mapper (Z z) map(tofrom: z.z)
+
+#pragma omp target data map(dummy)
+  {
+  #pragma omp target
+    {
+      yy.z++;
+      zz.z++;
+    }
+    yy.z++;
+  }
+
+  #pragma omp declare mapper(P x) map(to:x.zp) map(tofrom:*x.zp)
+
+  #pragma omp target
+  {
+    zz = *pp.zp;
+  }
+
+  return zz.z;
+}
+
+// { dg-final { scan-tree-dump-times {mapper_binding\(struct Y,omp declare mapper ~1Y\) mapper_binding\(struct Z,omp declare mapper ~1Z\)} 2 "original" } }
+// { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" } }
diff --git a/gcc/testsuite/g++.dg/gomp/ind-base-3.C b/gcc/testsuite/g++.dg/gomp/ind-base-3.C
index dbabaf7680c..7695b1f907e 100644
--- a/gcc/testsuite/g++.dg/gomp/ind-base-3.C
+++ b/gcc/testsuite/g++.dg/gomp/ind-base-3.C
@@ -24,7 +24,6 @@ int main (int argc, char *argv[])
     {
 #pragma omp target map(choose(&a, &b, i)->x[:10])
 /* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)->S::x\[0\]'} "" { target *-*-* } .-1 } */
-/* { dg-message {sorry, unimplemented: unsupported map expression 'choose\(\(& a\), \(& b\), i\)'} "" { target *-*-* } .-2 } */
       for (int j = 0; j < 10; j++)
         choose (&a, &b, i)->x[j]++;
     }
diff --git a/gcc/testsuite/g++.dg/gomp/member-array-2.C b/gcc/testsuite/g++.dg/gomp/member-array-2.C
index e60bb5585a1..caf8ece4262 100644
--- a/gcc/testsuite/g++.dg/gomp/member-array-2.C
+++ b/gcc/testsuite/g++.dg/gomp/member-array-2.C
@@ -70,7 +70,6 @@ main (int argc, char *argv[])
      it for now.  */
   #pragma omp target map(c.get_arr()[:100])
   /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)\[0\]'} "" { target *-*-* } .-1 } */
-  /* { dg-message {sorry, unimplemented: unsupported map expression 'c\.C::get_arr\(\)'} "" { target *-*-* } .-2 } */
   #pragma omp teams distribute parallel for
     for (int i = 0; i < 100; i++)
       c.get_arr()[i] += 2;
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 86a07c282af..c13564b85f5 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -345,6 +345,10 @@ enum omp_clause_code {
   /* OpenMP clause: has_device_addr (variable-list).  */
   OMP_CLAUSE_HAS_DEVICE_ADDR,
 
+  /* OpenMP mapper binding: record implicit mappers in scope for aggregate
+     types used within an offload region.  */
+  OMP_CLAUSE__MAPPER_BINDING_,
+
   /* Internal structure to hold OpenACC cache directive's variable-list.
      #pragma acc cache (variable-list).  */
   OMP_CLAUSE__CACHE_,
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index ccbafa98699..beb0073800c 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -973,6 +973,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	  pp_string (pp, "attach_zero_length_array_section");
 	  break;
+	case GOMP_MAP_UNSET:
+	  pp_string (pp, "unset");
+	  break;
+	case GOMP_MAP_PUSH_MAPPER_NAME:
+	  pp_string (pp, "push_mapper");
+	  break;
+	case GOMP_MAP_POP_MAPPER_NAME:
+	  pp_string (pp, "pop_mapper");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -1036,6 +1045,23 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 			 spc, flags, false);
       goto print_clause_size;
 
+    case OMP_CLAUSE__MAPPER_BINDING_:
+      pp_string (pp, "mapper_binding(");
+      if (OMP_CLAUSE__MAPPER_BINDING__ID (clause))
+	{
+	  dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__ID (clause), spc,
+			     flags, false);
+	  pp_comma (pp);
+	}
+      dump_generic_node (pp,
+			 TREE_TYPE (OMP_CLAUSE__MAPPER_BINDING__DECL (clause)),
+			 spc, flags, false);
+      pp_comma (pp);
+      dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__MAPPER (clause), spc,
+			 flags, false);
+      pp_right_paren (pp);
+      break;
+
     case OMP_CLAUSE_NUM_TEAMS:
       pp_string (pp, "num_teams(");
       if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause))
@@ -3806,6 +3832,22 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
       is_expr = false;
       break;
 
+    case OMP_DECLARE_MAPPER:
+      pp_string (pp, "#pragma omp declare mapper (");
+      if (OMP_DECLARE_MAPPER_ID (node))
+	{
+	  dump_generic_node (pp, OMP_DECLARE_MAPPER_ID (node), spc, flags,
+			     false);
+	  pp_colon (pp);
+	}
+      dump_generic_node (pp, OMP_DECLARE_MAPPER_TYPE (node), spc, flags, false);
+      pp_space (pp);
+      dump_generic_node (pp, OMP_DECLARE_MAPPER_DECL (node), spc, flags, false);
+      pp_right_paren (pp);
+      dump_omp_clauses (pp, OMP_DECLARE_MAPPER_CLAUSES (node), spc, flags);
+      is_expr = false;
+      break;
+
     case TRANSACTION_EXPR:
       if (TRANSACTION_EXPR_OUTER (node))
 	pp_string (pp, "__transaction_atomic [[outer]]");
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 007c9325b17..f28570d5e16 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -294,6 +294,7 @@ unsigned const char omp_clause_num_ops[] =
   2, /* OMP_CLAUSE_TO  */
   2, /* OMP_CLAUSE_MAP  */
   1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
+  3, /* OMP_CLAUSE__MAPPER_BINDING_  */
   2, /* OMP_CLAUSE__CACHE_  */
   2, /* OMP_CLAUSE_GANG  */
   1, /* OMP_CLAUSE_ASYNC  */
@@ -384,6 +385,7 @@ const char * const omp_clause_code_name[] =
   "to",
   "map",
   "has_device_addr",
+  "_mapper_binding_",
   "_cache_",
   "gang",
   "async",
diff --git a/gcc/tree.def b/gcc/tree.def
index f015021e9dc..d093927a0b5 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1247,6 +1247,13 @@ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1)
    Operand 0: OMP_MASTER_BODY: Master section body.  */
 DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1)
 
+/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN]
+   Operand 0: Identifier.
+   Operand 1: Type.
+   Operand 2: Variable decl.
+   Operand 3: List of clauses.  */
+DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 4)
+
 /* OpenACC - #pragma acc cache (variable1 ... variableN)
    Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
 	OMP_CLAUSE__CACHE_ clauses).  */
diff --git a/gcc/tree.h b/gcc/tree.h
index 31c7a129f3c..d580e9bcdc1 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1532,6 +1532,15 @@ class auto_suppress_location_wrappers
 #define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\
   TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0)
 
+#define OMP_DECLARE_MAPPER_ID(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 0)
+#define OMP_DECLARE_MAPPER_TYPE(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1)
+#define OMP_DECLARE_MAPPER_DECL(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2)
+#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 3)
+
 #define OMP_SCAN_BODY(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
 #define OMP_SCAN_CLAUSES(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
 
@@ -1960,6 +1969,18 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE__SCANTEMP__CONTROL(NODE) \
   TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__SCANTEMP_))
 
+#define OMP_CLAUSE__MAPPER_BINDING__ID(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+			OMP_CLAUSE__MAPPER_BINDING_), 0)
+
+#define OMP_CLAUSE__MAPPER_BINDING__DECL(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+			OMP_CLAUSE__MAPPER_BINDING_), 1)
+
+#define OMP_CLAUSE__MAPPER_BINDING__MAPPER(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+			OMP_CLAUSE__MAPPER_BINDING_), 2)
+
 /* SSA_NAME accessors.  */
 
 /* Whether SSA_NAME NODE is a virtual operand.  This simply caches the
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 84316f953d0..df706830551 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -182,7 +182,13 @@ enum gomp_map_kind
     /* An attach or detach operation.  Rewritten to the appropriate type during
        gimplification, depending on directive (i.e. "enter data" or
        parallel/kernels region vs. "exit data").  */
-    GOMP_MAP_ATTACH_DETACH =		(GOMP_MAP_LAST | 3)
+    GOMP_MAP_ATTACH_DETACH =		(GOMP_MAP_LAST | 3),
+    /* Unset, used for "declare mapper" maps with no explicit data movement
+       specified.  These use the movement specified at the invocation site.  */
+    GOMP_MAP_UNSET =			(GOMP_MAP_LAST | 4),
+    /* Used to record the name of a named mapper.  */
+    GOMP_MAP_PUSH_MAPPER_NAME =		(GOMP_MAP_LAST | 5),
+    GOMP_MAP_POP_MAPPER_NAME =		(GOMP_MAP_LAST | 6)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-1.C b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C
new file mode 100644
index 00000000000..aba4f426539
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-1.C
@@ -0,0 +1,87 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+#define N 64
+
+struct points
+{
+  double *x;
+  double *y;
+  double *z;
+  size_t len;
+};
+
+#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \
+				     map(p.x[0:p.len]) \
+				     map(p.y[0:p.len]) \
+				     map(p.z[0:p.len])
+
+struct shape
+{
+  points tmp;
+  points *pts;
+  int metadata[128];
+};
+
+#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp)
+
+void
+alloc_points (points *pts, size_t sz)
+{
+  pts->x = new double[sz];
+  pts->y = new double[sz];
+  pts->z = new double[sz];
+  pts->len = sz;
+  for (int i = 0; i < sz; i++)
+    pts->x[i] = pts->y[i] = pts->z[i] = 0;
+}
+
+int main (int argc, char *argv[])
+{
+  shape myshape;
+  points mypts;
+
+  myshape.pts = &mypts;
+
+  alloc_points (&myshape.tmp, N);
+  myshape.pts = new points;
+  alloc_points (myshape.pts, N);
+
+  #pragma omp target map(myshape)
+  {
+    for (int i = 0; i < N; i++)
+      {
+	myshape.pts->x[i]++;
+	myshape.pts->y[i]++;
+	myshape.pts->z[i]++;
+      }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (myshape.pts->x[i] == 1);
+      assert (myshape.pts->y[i] == 1);
+      assert (myshape.pts->z[i] == 1);
+    }
+
+  #pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      {
+	myshape.pts->x[i]++;
+	myshape.pts->y[i]++;
+	myshape.pts->z[i]++;
+      }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (myshape.pts->x[i] == 2);
+      assert (myshape.pts->y[i] == 2);
+      assert (myshape.pts->z[i] == 2);
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-2.C b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C
new file mode 100644
index 00000000000..d848fdb7369
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-2.C
@@ -0,0 +1,55 @@
+// { dg-do run }
+
+#include <cassert>
+
+#define N 256
+
+struct doublebuf
+{
+  int buf_a[N][N];
+  int buf_b[N][N];
+};
+
+#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \
+					   map(b.buf_b[0:N/2][0:N])
+
+#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \
+					   map(b.buf_b[N/2:N/2][0:N])
+
+int main (int argc, char *argv[])
+{
+  doublebuf db;
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      db.buf_a[i][j] = db.buf_b[i][j] = 0;
+
+  #pragma omp target map(mapper(lo), tofrom:db)
+  {
+    for (int i = 0; i < N / 2; i++)
+      for (int j = 0; j < N; j++)
+	{
+	  db.buf_a[i][j]++;
+	  db.buf_b[i][j]++;
+	}
+  }
+
+  #pragma omp target map(mapper(hi), tofrom:db)
+  {
+    for (int i = N / 2; i < N; i++)
+      for (int j = 0; j < N; j++)
+	{
+	  db.buf_a[i][j]++;
+	  db.buf_b[i][j]++;
+	}
+  }
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      {
+	assert (db.buf_a[i][j] == 1);
+	assert (db.buf_b[i][j] == 1);
+      }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-3.C b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C
new file mode 100644
index 00000000000..ea9b7ded75b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-3.C
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+  int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+  void incr_a (S my_s)
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 100; i++)
+	my_s.myarr[i]++;
+    }
+  }
+}
+
+namespace B
+{
+  void incr_b (S my_s)
+  {
+#pragma omp target
+    {
+      for (int i = 100; i < 200; i++)
+	my_s.myarr[i]++;
+    }
+  }
+}
+
+int main (int argc, char *argv[])
+{
+  S my_s;
+
+  my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+  {
+    for (int i = 0; i < 20; i++)
+      my_s.myarr[i]++;
+  }
+
+  A::incr_a (my_s);
+  B::incr_b (my_s);
+
+  for (int i = 0; i < 200; i++)
+    assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-4.C b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C
new file mode 100644
index 00000000000..f194e63b5b7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-4.C
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+  int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+  void incr_a (S &my_s)
+  {
+#pragma omp target
+    {
+      for (int i = 0; i < 100; i++)
+	my_s.myarr[i]++;
+    }
+  }
+}
+
+namespace B
+{
+  void incr_b (S &my_s)
+  {
+#pragma omp target
+    {
+      for (int i = 100; i < 200; i++)
+	my_s.myarr[i]++;
+    }
+  }
+}
+
+int main (int argc, char *argv[])
+{
+  S my_s;
+
+  my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+  {
+    for (int i = 0; i < 20; i++)
+      my_s.myarr[i]++;
+  }
+
+  A::incr_a (my_s);
+  B::incr_b (my_s);
+
+  for (int i = 0; i < 200; i++)
+    assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-5.C b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C
new file mode 100644
index 00000000000..0030de8791a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-5.C
@@ -0,0 +1,52 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+  int *myarr;
+  int len;
+};
+
+class C
+{
+  S smemb;
+#pragma omp declare mapper (custom:S s) map(to:s.myarr) \
+					map(tofrom:s.myarr[0:s.len])
+
+public:
+  C(int l)
+  {
+    smemb.myarr = new int[l];
+    smemb.len = l;
+    for (int i = 0; i < l; i++)
+      smemb.myarr[i] = 0;
+  }
+  void bump();
+  void check();
+};
+
+void
+C::bump ()
+{
+#pragma omp target map(mapper(custom), tofrom: smemb)
+  {
+    for (int i = 0; i < smemb.len; i++)
+      smemb.myarr[i]++;
+  }
+}
+
+void
+C::check ()
+{
+  for (int i = 0; i < smemb.len; i++)
+    assert (smemb.myarr[i] == 1);
+}
+
+int main (int argc, char *argv[])
+{
+  C test (100);
+  test.bump ();
+  test.check ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-6.C b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C
new file mode 100644
index 00000000000..14ed10df702
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-6.C
@@ -0,0 +1,37 @@
+// { dg-do run }
+
+#include <cassert>
+
+template <typename T>
+void adjust (T param)
+{
+#pragma omp declare mapper (T x) map(to:x.len, x.base) \
+				 map(tofrom:x.base[0:x.len])
+
+#pragma omp target
+  for (int i = 0; i < param.len; i++)
+    param.base[i]++;
+}
+
+struct S {
+  int len;
+  int *base;
+};
+
+int main (int argc, char *argv[])
+{
+  S a;
+
+  a.len = 100;
+  a.base = new int[a.len];
+
+  for (int i = 0; i < a.len; i++)
+    a.base[i] = 0;
+
+  adjust (a);
+
+  for (int i = 0; i < a.len; i++)
+    assert (a.base[i] == 1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-7.C b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C
new file mode 100644
index 00000000000..ab632099714
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-7.C
@@ -0,0 +1,48 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+  int *myarr;
+};
+
+struct T
+{
+  S *s;
+};
+
+#pragma omp declare mapper (s100: S x) map(to: x.myarr) \
+				       map(tofrom: x.myarr[0:100])
+
+void
+bump (T t)
+{
+  /* Here we have an implicit/default mapper invoking a named mapper.  We
+     need to make sure that can be located properly at gimplification
+     time.  */
+#pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0])
+
+#pragma omp target
+  for (int i = 0; i < 100; i++)
+    t.s->myarr[i]++;
+}
+
+int main (int argc, char *argv[])
+{
+  S my_s;
+  T my_t;
+
+  my_s.myarr = new int[100];
+  my_t.s = &my_s;
+
+  for (int i = 0; i < 100; i++)
+    my_s.myarr[i] = 0;
+
+  bump (my_t);
+
+  for (int i = 0; i < 100; i++)
+    assert (my_s.myarr[i] == 1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/declare-mapper-8.C b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C
new file mode 100644
index 00000000000..3818e5264d3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare-mapper-8.C
@@ -0,0 +1,61 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+  int *myarr;
+  int len;
+};
+
+template<typename T>
+class C
+{
+  T memb;
+#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \
+				 map(tofrom:t.myarr[0:t.len])
+
+public:
+  C(int sz);
+  ~C();
+  void bump();
+  void check();
+};
+
+template<typename T>
+C<T>::C(int sz)
+{
+  memb.myarr = new int[sz];
+  for (int i = 0; i < sz; i++)
+    memb.myarr[i] = 0;
+  memb.len = sz;
+}
+
+template<typename T>
+C<T>::~C()
+{
+  delete[] memb.myarr;
+}
+
+template<typename T>
+void C<T>::bump()
+{
+#pragma omp target map(memb)
+  for (int i = 0; i < memb.len; i++)
+    memb.myarr[i]++;
+}
+
+template<typename T>
+void C<T>::check()
+{
+  for (int i = 0; i < memb.len; i++)
+    assert (memb.myarr[i] == 1);
+}
+
+int main(int argc, char *argv[])
+{
+  C<S> c_int(100);
+  c_int.bump();
+  c_int.check();
+  return 0;
+}
