diff mbox series

[8/8] OpenMP 5.0 "declare mapper" support for C++

Message ID d8b937483709adb1b937a7bb5b1248e05fbe4455.1645214027.git.julian@codesourcery.com
State New
Headers show
Series OpenMP 5.0: C++ "declare mapper" support (plus struct rework, etc.) | expand

Commit Message

Julian Brown Feb. 18, 2022, 8:04 p.m. UTC
This patch implements OpenMP 5.0 "declare mapper" support for C++ --
except for arrays of structs with mappers, which are TBD. I've taken cues
from the existing "declare reduction" support where appropriate, though
obviously the details of implementation differ somewhat (in particular,
"declare mappers" must survive longer, until gimplification time).

Both named and unnamed (default) mappers are supported, and both
explicitly-mapped structures and implicitly-mapped struct-typed variables
used within an offload region are supported. For the latter, we need a
way to communicate to the middle end which mapper (visible, in-scope) is
the right one to use -- for that, we scan the target body in the front
end for uses of structure (etc.) types, and create artificial "mapper
binding" clauses to associate types with visible mappers. (It doesn't
matter too much if we create these mapper bindings a bit over-eagerly,
since they will only be used if needed later during gimplification.)

Another difficulty concerns the order in which an OpenMP offload region
body's clauses are processed relative to its body: in order to add
clauses for instantiated mappers, we need to have processed the body
already in order to find which variables have been used, but at present
the region's clauses are processed strictly before the body. So, this
patch adds a second clause-processing step after gimplification of the
body in order to add clauses resulting from instantiated mappers.

(C and Fortran FE support are TBD.)

Tested (alongside previous patches) with offloading to NVPTX, and
bootstrapped.  Any comments on this general approach?  Thoughts on
handling arrays of custom-mapped structs?  OK (for stage 1)?

Thanks,

Julian

2022-02-18  Julian Brown  <julian@codesourcery.com>

gcc/cp/
	* cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function.
	* cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES): Add new
	langhook.
	* cp-tree.h (lang_decl_fn): Add omp_declare_mapper_p field.
	(DECL_OMP_DECLARE_MAPPER_P): New function.
	(omp_mapper_id, cp_check_omp_declare_mapper, omp_instantiate_mappers,
	cxx_omp_finish_mapper_clauses): Add prototypes.
	* decl.cc (duplicate_decls): Support DECL_OMP_DECLARE_MAPPER_P
	functions.
	(finish_function): Likewise.
	* mangle.cc (decl_mangling_context): Likewise.
	* name-lookup.cc (set_decl_context_in_fn): Likewise.
	* parser.cc (cp_parser_class_specifier_1): Likewise.
	(cp_parser_omp_declare_mapper, cp_parser_omp_declare_mapper_maplist):
	New functions.
	(cp_parser_late_parsing_for_member): Support DECL_OMP_DECLARE_MAPPER_P
	functions.
	(cp_parser_omp_clause_map): Add KIND parameter.  Support "mapper"
	modifier.
	(cp_parser_omp_all_clauses): Add KIND argument to
	cp_parser_omp_clause_map call.
	(cp_parser_omp_target): Call omp_instantiate_mappers before
	finish_omp_clauses.
	(cp_parser_omp_declare): Add "declare mapper" support.
	* pt.cc (instantiate_class_template_1, tsubst_function_decl): Support
	DECL_OMP_DECLARE_MAPPER_P functions.
	(tsubst_omp_clauses): Call omp_instantiate_mappers before
	finish_omp_clauses, for target regions.
	(tsubst_expr): Support DECL_OMP_DECLARE_MAPPER_P functions.
	(tsubst_omp_udm): New function.
	(instantiate_body): Support DECL_OMP_DECLARE_MAPPER_P functions.
	* semantics.cc (gimplify.h): Include.
	(expand_or_defer_fn_1): Support DECL_OMP_DECLARE_MAPPER_P functions.
	(omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive,
	cp_check_omp_declare_mapper): New functions.
	(remap_mapper_decl_info): New struct.
	(remap_mapper_decl_1, omp_instantiate_mapper, omp_instantiate_mappers):
	New functions.
	(finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and
	GOMP_MAP_POP_MAPPER_NAME artificial clauses.
	(mapper_list): New struct.
	(find_nested_mappers): New function.
	(omp_target_walk_data): Add MAPPERS field.
	(finish_omp_target_clauses_r): Scan for uses of struct/union/class type
	variables.
	(finish_omp_target_clauses): Create artificial mapper binding clauses
	for used structs/unions/classes in offload region.

gcc/fortran/
	* parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes
	(for additions to omp-general.h).

gcc/
	* gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field.
	(new_omp_context): Initialise IMPLICIT_MAPPERS hash map.
	(delete_omp_context): Delete IMPLICIT_MAPPERS hash map.
	(instantiate_mapper_info, remap_mapper_decl_info): New structs.
	(remap_mapper_decl_1, omp_instantiate_mapper, handled_struct_p,
	omp_instantiate_implicit_mappers, new_omp_context_for_scan): New
	functions.
	(gimplify_scan_omp_clauses): Add optional USE_MAPPERS parameter.
	Instantiate implicit mappers if true.  Handle artificial mapper_binding
	clauses.
	(gimplify_omp_workshare): Update call to gimplify_scan_omp_clauses, and
	call again after gimplifying target body.
	(gimplify_omp_declare_mapper): New function.
	(gimplify_expr): Call above function.
	* langhooks.def (lhd_omp_finish_mapper_clauses): Add prototype.
	(LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES): Define.
	(LANG_HOOK_DECLS): Add above macro.
	* langhooks.cc (lhd_omp_finish_mapper_clauses): New function.
	* langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES
	hook.
	* omp-general.h (omp_name_type): Add struct, hash type traits.
	* tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_.
	* tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET,
	GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
	clauses.  Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER.
	* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add
	OMP_CLAUSE__MAPPER_BINDING_.
	* tree.def (OMP_DECLARE_MAPPER): New tree code.
	* tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_TYPE,
	OMP_DECLARE_MAPPER_DECL, OMP_DECLARE_MAPPER_CLAUSES): New defines.
	(OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL,
	OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines.

include/
	* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET,
	GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
	clause types.

gcc/testsuite/
	* c-c++-common/gomp/map-6.c: Update error scan output.
	* g++.dg/gomp/declare-mapper-1.C: New test.
	* g++.dg/gomp/declare-mapper-2.C: New test.
	* g++.dg/gomp/declare-mapper-3.C: New test.
	* g++.dg/gomp/declare-mapper-4.C: New test.

libgomp/
	* testsuite/libgomp.c++/declare-mapper-1.C: New test.
	* testsuite/libgomp.c++/declare-mapper-2.C: New test.
	* testsuite/libgomp.c++/declare-mapper-3.C: New test.
	* testsuite/libgomp.c++/declare-mapper-4.C: New test.
	* testsuite/libgomp.c++/declare-mapper-5.C: New test.
	* testsuite/libgomp.c++/declare-mapper-6.C: New test.
	* testsuite/libgomp.c++/declare-mapper-7.C: New test.
	* testsuite/libgomp.c++/declare-mapper-8.C: New test.
---
 gcc/cp/cp-gimplify.cc                         |   6 +
 gcc/cp/cp-objcp-common.h                      |   2 +
 gcc/cp/cp-tree.h                              |  10 +
 gcc/cp/decl.cc                                |  18 +-
 gcc/cp/mangle.cc                              |   5 +-
 gcc/cp/name-lookup.cc                         |   3 +-
 gcc/cp/parser.cc                              | 370 ++++++++++++-
 gcc/cp/pt.cc                                  |  92 +++-
 gcc/cp/semantics.cc                           | 496 +++++++++++++++++-
 gcc/fortran/parse.cc                          |   3 +
 gcc/gimplify.cc                               | 316 ++++++++++-
 gcc/langhooks-def.h                           |   3 +
 gcc/langhooks.cc                              |   9 +
 gcc/langhooks.h                               |   4 +
 gcc/omp-general.h                             |  52 ++
 gcc/testsuite/c-c++-common/gomp/map-6.c       |   6 +-
 gcc/testsuite/g++.dg/gomp/declare-mapper-1.C  |  58 ++
 gcc/testsuite/g++.dg/gomp/declare-mapper-2.C  |  30 ++
 gcc/testsuite/g++.dg/gomp/declare-mapper-3.C  |  27 +
 gcc/testsuite/g++.dg/gomp/declare-mapper-4.C  |  74 +++
 gcc/tree-core.h                               |   4 +
 gcc/tree-pretty-print.cc                      |  42 ++
 gcc/tree.cc                                   |   2 +
 gcc/tree.def                                  |   7 +
 gcc/tree.h                                    |  21 +
 include/gomp-constants.h                      |   8 +-
 .../testsuite/libgomp.c++/declare-mapper-1.C  |  87 +++
 .../testsuite/libgomp.c++/declare-mapper-2.C  |  55 ++
 .../testsuite/libgomp.c++/declare-mapper-3.C  |  63 +++
 .../testsuite/libgomp.c++/declare-mapper-4.C  |  63 +++
 .../testsuite/libgomp.c++/declare-mapper-5.C  |  52 ++
 .../testsuite/libgomp.c++/declare-mapper-6.C  |  37 ++
 .../testsuite/libgomp.c++/declare-mapper-7.C  |  48 ++
 .../testsuite/libgomp.c++/declare-mapper-8.C  |  61 +++
 34 files changed, 2090 insertions(+), 44 deletions(-)
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-2.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-3.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-4.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-1.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-3.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-4.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-5.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-6.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-7.C
 create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-8.C
diff mbox series

Patch

diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc
index d7323fb5c09..e5d59766e5c 100644
--- a/gcc/cp/cp-gimplify.cc
+++ b/gcc/cp/cp-gimplify.cc
@@ -2284,6 +2284,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 4c137313525..6a0df9cc913 100644
--- a/gcc/cp/cp-objcp-common.h
+++ b/gcc/cp/cp-objcp-common.h
@@ -182,6 +182,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_MAPPABLE_TYPE
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index f253b32c3f2..8f634197dcc 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -2891,6 +2891,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;
@@ -4295,6 +4296,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) \
@@ -7653,10 +7659,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> &);
@@ -8208,6 +8217,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 7b48b56231b..bc3196d8c12 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -1859,6 +1859,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)
@@ -17821,7 +17833,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
@@ -17898,7 +17911,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 a20f0e00329..4aba371f262 100644
--- a/gcc/cp/mangle.cc
+++ b/gcc/cp/mangle.cc
@@ -931,10 +931,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 93c4eb7193b..b09d6c142b6 100644
--- a/gcc/cp/name-lookup.cc
+++ b/gcc/cp/name-lookup.cc
@@ -3354,7 +3354,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 8d5ae9c44d0..e914a314934 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -26203,10 +26203,12 @@  cp_parser_class_specifier_1 (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
@@ -32133,6 +32135,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.  */
@@ -32193,6 +32197,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,
@@ -39488,10 +39502,9 @@  cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
      always | close */
 
 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;
@@ -39509,11 +39522,24 @@  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)->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);
@@ -39536,6 +39562,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)
 	{
@@ -39549,20 +39576,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)
@@ -39603,8 +39693,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;
 }
@@ -40405,7 +40517,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:
@@ -44549,6 +44661,7 @@  cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
+  clauses = omp_instantiate_mappers (clauses);
   clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
 
   c_omp_adjust_map_clauses (clauses, true);
@@ -46491,6 +46604,235 @@  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)
+	return false;
+      maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+    }
+
+  if (maplist == NULL_TREE)
+    {
+      cp_parser_error (parser, "missing %<map%> clause");
+      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))
+    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) \
@@ -46531,6 +46873,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);
@@ -46544,7 +46892,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 ca763f1a805..f09248b09f1 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -12093,9 +12093,13 @@  instantiate_class_template_1 (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)))
@@ -14129,6 +14133,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
@@ -17852,6 +17864,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))
@@ -18568,6 +18582,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
 		  {
@@ -19513,6 +19531,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.  */
 
@@ -26274,7 +26352,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)
@@ -26373,6 +26452,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,
@@ -26400,6 +26482,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 f1214de0906..2a188993763 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
@@ -4729,7 +4730,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);
@@ -4742,6 +4744,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)))))
@@ -4775,7 +4778,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;
@@ -5959,6 +5963,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).
@@ -6440,6 +6514,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.  */
@@ -6700,6 +6799,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.  */
 
@@ -7942,6 +8277,13 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE__CACHE_:
 	  t = OMP_CLAUSE_DECL (c);
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME))
+	    {
+	      remove = true;
+	      break;
+	    }
 	  if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	    {
 	      if (handle_omp_array_sections (c, ort))
@@ -9286,6 +9628,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
@@ -9311,6 +9755,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
@@ -9324,6 +9770,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.  */
@@ -9337,6 +9785,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);
@@ -9440,10 +9904,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 db918291b10..3caa4d4b4b5 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 12d508c764c..b0fb15084e8 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.  */
@@ -451,6 +452,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;
@@ -474,6 +476,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);
 }
@@ -10350,23 +10353,176 @@  error_out:
   return success;
 }
 
-/* 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 instantiate_mapper_info
 {
-  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;
+  vec<omp_mapping_group> *groups;
+  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+  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;
+}
+
+static bool
+handled_struct_p (vec<omp_mapping_group> *groups, tree s)
+{
+  if (!groups)
+    return false;
+
+  for (auto &i : *groups)
+    {
+      tree node = *i.grp_start;
+      if (OMP_CLAUSE_CODE (node) == OMP_CLAUSE_MAP
+	  && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_STRUCT
+	  && OMP_CLAUSE_DECL (node) == s)
+	return true;
+    }
+  return false;
+}
+
+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 = handled_struct_p (im_info->groups, 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 ())
@@ -10390,6 +10546,34 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       default:
 	break;
       }
+  gimplify_omp_ctxp = ctx;
+  return ctx;
+}
+
+/* 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 use_mappers = 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;
+  tree mapper_clauses = NULL_TREE;
+
+  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 (code == OMP_TARGET
       || code == OMP_TARGET_DATA
@@ -10398,11 +10582,50 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
     {
       vec<omp_mapping_group> *groups;
       groups = omp_gather_mapping_groups (list_p);
+      hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+
+      if (groups)
+	grpmap = omp_index_mapping_groups (groups);
+
+      if (use_mappers)
+	{
+	  instantiate_mapper_info im_info;
+
+	  im_info.groups = groups;
+	  im_info.grpmap = grpmap;
+	  im_info.mapper_clauses_p = &mapper_clauses;
+	  im_info.omp_ctx = ctx;
+
+	  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);
+
+	  if (groups)
+	    {
+	      delete grpmap;
+	      delete groups;
+	      grpmap = NULL;
+	    }
+
+	  if (!mapper_clauses)
+	    {
+	      gimplify_omp_ctxp = ctx;
+	      return;
+	    }
+
+	  list_p = &mapper_clauses;
+	  groups = omp_gather_mapping_groups (list_p);
+
+	  if (groups)
+	    grpmap = omp_index_mapping_groups (groups);
+	}
+
       if (groups)
 	{
-	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
-	  grpmap = omp_index_mapping_groups (groups);
-
 	  omp_build_struct_sibling_lists (code, region_type, groups, &grpmap);
 
 	  omp_mapping_group *outlist = NULL;
@@ -10445,6 +10668,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	}
     }
 
+  gcc_assert (!use_mappers || mapper_clauses);
+
   while ((c = *list_p) != NULL)
     {
       bool remove = false;
@@ -11200,6 +11425,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;
@@ -11739,7 +11987,15 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
-  ctx->clauses = *orig_list_p;
+  if (mapper_clauses)
+    {
+      /* Put clauses implicitly created from "declare mapper" mappers
+	 at the front of the clause list.  */
+      *list_p = *orig_list_p;
+      *orig_list_p = ctx->clauses = mapper_clauses;
+    }
+  else
+    ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
 }
 
@@ -14913,8 +15169,10 @@  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;
+  struct gimplify_omp_ctx *ctx
+    = new_omp_context_for_scan (ort, TREE_CODE (expr));
   gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
-			     TREE_CODE (expr));
+			     TREE_CODE (expr), ctx, false);
   if (TREE_CODE (expr) == OMP_TARGET)
     optimize_target_teams (expr, pre_p);
   if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -14952,6 +15210,12 @@  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)
+    gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+			       TREE_CODE (expr), ctx, true);
   gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr),
 			       TREE_CODE (expr));
   in_omp_construct = save_in_omp_construct;
@@ -15611,6 +15875,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
@@ -16520,6 +16794,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 49c8f5820cf..fa49092636a 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -84,6 +84,7 @@  extern enum omp_clause_default_kind lhd_omp_predetermined_sharing (tree);
 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_finish_mapper_clauses (tree);
 struct gimplify_omp_ctx;
 extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
@@ -270,6 +271,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
@@ -303,6 +305,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 df970678a08..fc51dbe720a 100644
--- a/gcc/langhooks.cc
+++ b/gcc/langhooks.cc
@@ -634,6 +634,15 @@  lhd_omp_finish_clause (tree, gimple_seq *, bool)
 {
 }
 
+/* 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 0eec1b0f7ad..3bdc12badc9 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -306,6 +306,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 c0cf5f014cd..f676cc7c493 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -149,4 +149,56 @@  get_openacc_privatization_dump_flags ()
   return l_dump_flags;
 }
 
+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 19db264e805..24fe31d903f 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..c9b0a488607
--- /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\(tofrom: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\(to: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/tree-core.h b/gcc/tree-core.h
index bf2efa61330..6f66d16be78 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -342,6 +342,10 @@  enum omp_clause_code {
      OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list).  */
   OMP_CLAUSE_MAP,
 
+  /* 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 86fc5c090d6..58336a9adcd 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -940,6 +940,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 ();
 	}
@@ -1003,6 +1012,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))
@@ -3769,6 +3795,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 dfcdf6822f1..09580a09d0a 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -289,6 +289,7 @@  unsigned const char omp_clause_num_ops[] =
   2, /* OMP_CLAUSE_FROM  */
   2, /* OMP_CLAUSE_TO  */
   2, /* OMP_CLAUSE_MAP  */
+  3, /* OMP_CLAUSE__MAPPER_BINDING_  */
   2, /* OMP_CLAUSE__CACHE_  */
   2, /* OMP_CLAUSE_GANG  */
   1, /* OMP_CLAUSE_ASYNC  */
@@ -378,6 +379,7 @@  const char * const omp_clause_code_name[] =
   "from",
   "to",
   "map",
+  "_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 95334b077da..5bf82271f3e 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1475,6 +1475,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)
 
@@ -1885,6 +1894,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 b5933a0cd55..6b476ef621f 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;
+}