diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 87e999becd5d..e10af915d7ca 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1257,6 +1257,9 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree);
 extern tree c_omp_check_context_selector (location_t, tree);
 extern void c_omp_mark_declare_variant (location_t, tree, tree);
 extern void c_omp_adjust_map_clauses (tree, bool);
+template<typename T> struct omp_mapper_list;
+extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree);
+extern tree c_omp_instantiate_mappers (tree);
 
 namespace omp_addr_tokenizer { struct omp_addr_token; }
 typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 4cade64f22b7..1e159fe484b8 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3937,6 +3937,306 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
   return error_mark_node;
 }
 
+/* Given a mapper function MAPPER_FN, recursively scan through the map clauses
+   for that mapper, and if any of those should use a (named or unnamed) mapper
+   themselves, add it to MLIST.  */
+
+void
+c_omp_find_nested_mappers (omp_mapper_list<tree> *mlist, tree mapper_fn)
+{
+  tree mapper = lang_hooks.decls.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);
+
+	  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
+	    = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type);
+
+	  if (nested_mapper_fn)
+	    {
+	      mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
+	      c_omp_find_nested_mappers (mlist, nested_mapper_fn);
+	    }
+	  else if (mapper_name)
+	    {
+	      error ("mapper %qE not found for type %qT", mapper_name,
+		     elem_type);
+	      continue;
+	    }
+	}
+    }
+}
+
+struct remap_mapper_decl_info
+{
+  tree dummy_var;
+  tree expr;
+};
+
+/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl.  */
+
+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;
+}
+
+/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
+   OUTLIST.  OUTER_KIND is the mapping kind to use if not already specified in
+   the mapper declaration.  */
+
+static tree *
+omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+			enum gomp_map_kind outer_kind)
+{
+  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)
+	{
+	  location_t loc = OMP_CLAUSE_LOCATION (c);
+	  tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+	  if (t2 == t)
+	    {
+	      nonunit_array_with_mapper = true;
+	      /* We'd want use the mapper for the element type if this worked:
+		 look that one up.  */
+	      type = TREE_TYPE (TREE_TYPE (t));
+	    }
+	  else
+	    {
+	      t = t2;
+	      type = TREE_TYPE (t);
+	    }
+	}
+      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 = lang_hooks.decls.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
+	    = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+	  if (nested_mapper != mapper)
+	    {
+	      if (clause_kind == GOMP_MAP_UNSET)
+		clause_kind = outer_kind;
+
+	      outlist = omp_instantiate_mapper (outlist, nested_mapper,
+						t, clause_kind);
+	      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;
+}
+
+/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper
+   appropriate to the type of the data in that clause, if such a mapper is
+   visible in the current parsing context.  */
+
+tree
+c_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;
+	      }
+
+	    if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+	      {
+		location_t loc = OMP_CLAUSE_LOCATION (c);
+		tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+		if (t2 == t)
+		  {
+		    /* !!! Array sections of size >1 with mappers for elements
+		       are hard to support.  Do something here.  */
+		    nonunit_array_with_mapper = true;
+		    type = TREE_TYPE (TREE_TYPE (t));
+		  }
+		else
+		  {
+		    t = t2;
+		    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
+	      = lang_hooks.decls.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
+		  = lang_hooks.decls.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;
+}
+
 const struct c_omp_directive c_omp_directives[] = {
   /* Keep this alphabetically sorted by the first word.  Non-null second/third
      if any should precede null ones.  */
diff --git a/gcc/cp/constexpr.cc b/gcc/cp/constexpr.cc
index a661096192e5..0369d8d859c4 100644
--- a/gcc/cp/constexpr.cc
+++ b/gcc/cp/constexpr.cc
@@ -3101,6 +3101,9 @@ reduced_constant_expression_p (tree t)
       /* Even if we can't lower this yet, it's constant.  */
       return true;
 
+    case OMP_DECLARE_MAPPER:
+      return true;
+
     case CONSTRUCTOR:
       /* And we need to handle PTRMEM_CST wrapped in a CONSTRUCTOR.  */
       tree field;
@@ -6895,6 +6898,7 @@ cxx_eval_constant_expression (const constexpr_ctx *ctx, tree t,
     case LABEL_EXPR:
     case CASE_LABEL_EXPR:
     case PREDICT_EXPR:
+    case OMP_DECLARE_MAPPER:
       return t;
 
     case PARM_DECL:
@@ -9276,6 +9280,11 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now,
 	error_at (loc, "expression %qE is not a constant expression", t);
       return false;
 
+    case OMP_DECLARE_MAPPER:
+      /* This can be used to initialize VAR_DECLs: it's treated as a magic
+	 constant.  */
+      return true;
+
     case ASM_EXPR:
       if (flags & tf_error)
 	inline_asm_in_constexpr_error (loc);
diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc
index 28c339869b83..d87ce9dd883d 100644
--- a/gcc/cp/cp-gimplify.cc
+++ b/gcc/cp/cp-gimplify.cc
@@ -2302,6 +2302,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 1a67f14d9b38..49c480c1d21b 100644
--- a/gcc/cp/cp-objcp-common.h
+++ b/gcc/cp/cp-objcp-common.h
@@ -185,6 +185,15 @@ 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_MAPPER_LOOKUP
+#define LANG_HOOKS_OMP_MAPPER_LOOKUP cxx_omp_mapper_lookup
+#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
+  cxx_omp_extract_mapper_directive
+#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION cxx_omp_map_array_section
 #undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
 #define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
 #undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 29319bde2ca7..7ded32f4660c 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -2860,7 +2860,10 @@ struct GTY(()) lang_decl_base {
   /* VAR_DECL or FUNCTION_DECL has keyed decls.     */
   unsigned module_keyed_decls_p : 1;
 
-  /* 12 spare bits.  */
+  /* VAR_DECL being used to represent an OpenMP declared mapper.  */
+  unsigned omp_declare_mapper_p : 1;
+
+  /* 10 spare bits.  */
 };
 
 /* True for DECL codes which have template info and access.  */
@@ -4334,6 +4337,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) \
+  (DECL_LANG_SPECIFIC (VAR_DECL_CHECK (NODE))->u.base.omp_declare_mapper_p)
+
 /* Nonzero if DECL has been declared threadprivate by
    #pragma omp threadprivate.  */
 #define CP_DECL_THREADPRIVATE_P(DECL) \
@@ -7708,10 +7716,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> &);
@@ -8271,6 +8282,10 @@ 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 tree cxx_omp_mapper_lookup		(tree, tree);
+extern tree cxx_omp_extract_mapper_directive	(tree);
+extern tree cxx_omp_map_array_section		(location_t, 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 85b892cddf04..ff7745bd1e31 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -7311,6 +7311,12 @@ check_initializer (tree decl, tree init, int flags, vec<tree, va_gc> **cleanups)
     }
   else if (!init && DECL_REALLY_EXTERN (decl))
     ;
+  else if (flag_openmp
+	   && VAR_P (decl)
+	   && DECL_LANG_SPECIFIC (decl)
+	   && DECL_OMP_DECLARE_MAPPER_P (decl)
+	   && TREE_CODE (init) == OMP_DECLARE_MAPPER)
+    return NULL_TREE;
   else if (init || type_build_ctor_call (type)
 	   || TYPE_REF_P (type))
     {
@@ -8454,14 +8460,23 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p,
 	  varpool_node::get_create (decl);
 	}
 
+      if (flag_openmp
+	  && VAR_P (decl)
+	  && DECL_LANG_SPECIFIC (decl)
+	  && DECL_OMP_DECLARE_MAPPER_P (decl)
+	  && init)
+	{
+	  gcc_assert (TREE_CODE (init) == OMP_DECLARE_MAPPER);
+	  DECL_INITIAL (decl) = init;
+	}
       /* Convert the initializer to the type of DECL, if we have not
 	 already initialized DECL.  */
-      if (!DECL_INITIALIZED_P (decl)
-	  /* If !DECL_EXTERNAL then DECL is being defined.  In the
-	     case of a static data member initialized inside the
-	     class-specifier, there can be an initializer even if DECL
-	     is *not* defined.  */
-	  && (!DECL_EXTERNAL (decl) || init))
+      else if (!DECL_INITIALIZED_P (decl)
+	       /* If !DECL_EXTERNAL then DECL is being defined.  In the
+		  case of a static data member initialized inside the
+		  class-specifier, there can be an initializer even if DECL
+		  is *not* defined.  */
+	       && (!DECL_EXTERNAL (decl) || init))
 	{
 	  cleanups = make_tree_vector ();
 	  init = check_initializer (decl, init, flags, &cleanups);
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index 21f591b38f7c..ee913d8fa984 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -5772,10 +5772,15 @@ mark_used (tree decl, tsubst_flags_t complain /* = tf_warning_or_error */)
 
   /* If DECL has a deduced return type, we need to instantiate it now to
      find out its type.  For OpenMP user defined reductions, we need them
-     instantiated for reduction clauses which inline them by hand directly.  */
+     instantiated for reduction clauses which inline them by hand directly.
+     OpenMP declared mappers are used implicitly so must be instantiated
+     before they can be detected.  */
   if (undeduced_auto_decl (decl)
       || (TREE_CODE (decl) == FUNCTION_DECL
-	  && DECL_OMP_DECLARE_REDUCTION_P (decl)))
+	  && DECL_OMP_DECLARE_REDUCTION_P (decl))
+      || (TREE_CODE (decl) == VAR_DECL
+	  && DECL_LANG_SPECIFIC (decl)
+	  && DECL_OMP_DECLARE_MAPPER_P (decl)))
     maybe_instantiate_decl (decl);
 
   if (processing_template_decl || in_template_function ())
diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc
index 000f3780ef27..74e8fef8300d 100644
--- a/gcc/cp/error.cc
+++ b/gcc/cp/error.cc
@@ -1136,12 +1136,37 @@ dump_global_iord (cxx_pretty_printer *pp, tree t)
   pp_printf (pp, p, DECL_SOURCE_FILE (t));
 }
 
+/* Write a representation of OpenMP "declare mapper" T to PP in a manner
+   suitable for error messages.  */
+
+static void
+dump_omp_declare_mapper (cxx_pretty_printer *pp, tree t, int flags)
+{
+  pp_string (pp, "#pragma omp declare mapper");
+  if (t == NULL_TREE || t == error_mark_node)
+    return;
+  pp_space (pp);
+  pp_cxx_left_paren (pp);
+  if (OMP_DECLARE_MAPPER_ID (t))
+    {
+      pp_cxx_tree_identifier (pp, OMP_DECLARE_MAPPER_ID (t));
+      pp_colon (pp);
+    }
+  dump_type (pp, TREE_TYPE (t), flags);
+  pp_cxx_right_paren (pp);
+}
+
 static void
 dump_simple_decl (cxx_pretty_printer *pp, tree t, tree type, int flags)
 {
   if (TREE_CODE (t) == VAR_DECL && DECL_NTTP_OBJECT_P (t))
     return dump_expr (pp, DECL_INITIAL (t), flags);
 
+  if (TREE_CODE (t) == VAR_DECL
+      && DECL_LANG_SPECIFIC (t)
+      && DECL_OMP_DECLARE_MAPPER_P (t))
+    return dump_omp_declare_mapper (pp, DECL_INITIAL (t), flags);
+
   if (flags & TFF_DECL_SPECIFIERS)
     {
       if (concept_definition_p (t))
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 2ccf4fc4f8a7..c55cad836af0 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -39987,13 +39987,12 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
    map ( [map-type-modifier[,] ...] map-kind: variable-list )
 
    map-type-modifier:
-     always | close */
+     always | close | mapper ( mapper-name )  */
 
 static tree
-cp_parser_omp_clause_map (cp_parser *parser, tree list)
+cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind)
 {
   tree nlist, c;
-  enum gomp_map_kind kind = GOMP_MAP_TOFROM;
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
@@ -40011,11 +40010,16 @@ 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)
+	pos = cp_parser_skip_balanced_tokens (parser, pos + 1);
       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);
@@ -40038,6 +40042,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)
 	{
@@ -40051,20 +40056,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)
@@ -40110,8 +40178,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 					  NULL, true);
   finish_scope ();
 
+  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;
 }
@@ -40927,7 +41017,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:
@@ -45095,6 +45185,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
+  if (!processing_template_decl)
+    clauses = c_omp_instantiate_mappers (clauses);
   clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
 
   c_omp_adjust_map_clauses (clauses, true);
@@ -47359,6 +47451,172 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
   obstack_free (&declarator_obstack, p);
 }
 
+/* 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;
+  tree type = NULL_TREE, vardecl = 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, id, placeholder, mapper, maplist = NULL_TREE;
+
+  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);
+
+  vardecl = build_lang_decl (VAR_DECL, mapper_id, type);
+  DECL_ARTIFICIAL (vardecl) = 1;
+  TREE_STATIC (vardecl) = 1;
+  TREE_PUBLIC (vardecl) = 0;
+  DECL_EXTERNAL (vardecl) = 0;
+  DECL_DECLARED_CONSTEXPR_P (vardecl) = 1;
+  DECL_INITIALIZED_BY_CONSTANT_EXPRESSION_P (vardecl) = 1;
+  DECL_OMP_DECLARE_MAPPER_P (vardecl) = 1;
+
+  keep_next_level (true);
+  block = begin_omp_structured_block ();
+
+  if (block_scope)
+    DECL_CONTEXT (vardecl) = current_function_decl;
+  else if (current_class_type)
+    DECL_CONTEXT (vardecl) = current_class_type;
+  else
+    DECL_CONTEXT (vardecl) = current_namespace;
+
+  if (processing_template_decl)
+    vardecl = push_template_decl (vardecl);
+
+  id = cp_parser_declarator_id (parser, false);
+
+  if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+    {
+      finish_omp_structured_block (block);
+      goto fail;
+    }
+
+  placeholder = build_lang_decl (VAR_DECL, id, type);
+  DECL_CONTEXT (placeholder) = DECL_CONTEXT (vardecl);
+  if (processing_template_decl)
+    placeholder = push_template_decl (placeholder);
+  pushdecl (placeholder);
+  cp_finish_decl (placeholder, NULL_TREE, 0, NULL_TREE, 0);
+  DECL_ARTIFICIAL (placeholder) = 1;
+  TREE_USED (placeholder) = 1;
+
+  while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+    {
+      pragma_omp_clause c_kind = cp_parser_omp_clause_name (parser);
+      if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+	{
+	  if (c_kind != PRAGMA_OMP_CLAUSE_NONE)
+	    cp_parser_error (parser, "unexpected clause");
+	  finish_omp_structured_block (block);
+	  goto fail;
+	}
+      maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+      if (maplist == NULL_TREE)
+	break;
+    }
+
+  if (maplist == NULL_TREE)
+    {
+      cp_parser_error (parser, "missing %<map%> clause");
+      finish_omp_structured_block (block);
+      goto fail;
+    }
+
+  mapper = make_node (OMP_DECLARE_MAPPER);
+  TREE_TYPE (mapper) = type;
+  OMP_DECLARE_MAPPER_ID (mapper) = mapper_name;
+  OMP_DECLARE_MAPPER_DECL (mapper) = placeholder;
+  OMP_DECLARE_MAPPER_CLAUSES (mapper) = maplist;
+
+  finish_omp_structured_block (block);
+
+  DECL_INITIAL (vardecl) = mapper;
+
+  if (current_class_type)
+    {
+      if (processing_template_decl)
+	{
+	  retrofit_lang_decl (vardecl);
+	  SET_DECL_VAR_DECLARED_INLINE_P (vardecl);
+	}
+      finish_static_data_member_decl (vardecl, mapper,
+				      /*init_const_expr_p=*/true, NULL_TREE, 0);
+      finish_member_declaration (vardecl);
+    }
+  else if (processing_template_decl && block_scope)
+    add_decl_expr (vardecl);
+  else
+    pushdecl (vardecl);
+
+  cp_check_omp_declare_mapper (vardecl);
+
+  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) \
@@ -47399,6 +47657,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);
@@ -47412,7 +47676,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 3b4a3c845a14..61f2231cba6b 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -15220,6 +15220,13 @@ tsubst_decl (tree t, tree args, tsubst_flags_t complain)
 						  TYPE_ALIGN (TREE_TYPE (t)));
 	  }
 
+	if (flag_openmp
+	    && VAR_P (t)
+	    && DECL_LANG_SPECIFIC (t)
+	    && DECL_OMP_DECLARE_MAPPER_P (t)
+	    && strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL)
+	  DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), TREE_TYPE (r));
+
 	layout_decl (r, 0);
       }
       break;
@@ -18095,6 +18102,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
   new_clauses = nreverse (new_clauses);
   if (ort != C_ORT_OMP_DECLARE_SIMD)
     {
+      if (ort == C_ORT_OMP_TARGET)
+	new_clauses = c_omp_instantiate_mappers (new_clauses);
       new_clauses = finish_omp_clauses (new_clauses, ort);
       if (linear_no_step)
 	for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -19639,6 +19648,22 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 	}
       break;
 
+    case OMP_DECLARE_MAPPER:
+      {
+	t = copy_node (t);
+
+	tree decl = OMP_DECLARE_MAPPER_DECL (t);
+	decl = tsubst (decl, args, complain, in_decl);
+	tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
+	tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t);
+	clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args,
+				      complain, in_decl);
+	TREE_TYPE (t) = type;
+	OMP_DECLARE_MAPPER_DECL (t) = decl;
+	OMP_DECLARE_MAPPER_CLAUSES (t) = clauses;
+	RETURN (t);
+      }
+
     case TRANSACTION_EXPR:
       {
 	int flags = 0;
@@ -26920,7 +26945,9 @@ instantiate_decl (tree d, bool defer_ok, bool expl_inst_class_mem_p)
       || (external_p && VAR_P (d))
       /* Handle here a deleted function too, avoid generating
 	 its body (c++/61080).  */
-      || deleted_p)
+      || deleted_p
+      /* We need the initializer for an OpenMP declare mapper.  */
+      || (VAR_P (d) && DECL_LANG_SPECIFIC (d) && DECL_OMP_DECLARE_MAPPER_P (d)))
     {
       /* The definition of the static data member is now required so
 	 we must substitute the initializer.  */
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 0357a3e58d4d..0d18ab8ca699 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
@@ -5944,6 +5945,98 @@ 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);
+}
+
+tree
+cxx_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);
+}
+
+tree
+cxx_omp_extract_mapper_directive (tree vardecl)
+{
+  gcc_assert (TREE_CODE (vardecl) == VAR_DECL);
+
+  /* Instantiate the decl if we haven't already.  */
+  mark_used (vardecl);
+  tree body = DECL_INITIAL (vardecl);
+
+  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;
+}
+
+/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
+   nothing more complicated.  */
+
+tree
+cxx_omp_map_array_section (location_t loc, tree t)
+{
+  tree low = TREE_OPERAND (t, 1);
+  tree len = TREE_OPERAND (t, 2);
+
+  if (len && integer_onep (len))
+    {
+      t = TREE_OPERAND (t, 0);
+
+      if (!low)
+	low = integer_zero_node;
+
+      if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
+	t = convert_from_reference (t);
+
+      t = build_array_ref (loc, t, low);
+    }
+
+  return t;
+}
+
 /* Helper function for cp_parser_omp_declare_reduction_exprs
    and tsubst_omp_udr.
    Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6425,6 +6518,29 @@ 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_TYPE (udm);
+  location_t loc = DECL_SOURCE_LOCATION (udm);
+
+  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.  */
@@ -7965,6 +8081,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_MAP:
 	  if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
 	    goto move_implicit;
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+	    {
+	      remove = true;
+	      break;
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -9414,6 +9536,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;
+
+  omp_mapper_list<tree> *mappers;
 };
 
 /* Helper function of finish_omp_target_clauses, called via
@@ -9427,6 +9551,7 @@ 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;
+  omp_mapper_list<tree> *mlist = data->mappers;
 
   /* References inside of these expression codes shouldn't incur any
      form of mapping, so return early.  */
@@ -9440,6 +9565,27 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
   if (TREE_CODE (t) == OMP_CLAUSE)
     return NULL_TREE;
 
+  if (!processing_template_decl)
+    {
+      tree aggr_type = 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 = cxx_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);
@@ -9542,10 +9688,48 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
   else
     data.current_closure = NULL_TREE;
 
-  cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
-
   auto_vec<tree, 16> new_clauses;
 
+  if (!processing_template_decl)
+    {
+      hash_set<omp_name_type<tree> > seen_types;
+      auto_vec<tree> mapper_fns;
+      omp_mapper_list<tree> 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)
+	c_omp_find_nested_mappers (&mlist, mapper_fn);
+
+      FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
+	{
+	  tree mapper = cxx_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);
+	}
+    }
+  else
+    {
+      data.mappers = NULL;
+      cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r,
+				       &data);
+    }
+
   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 f04fd13cc696..b50f578af8f1 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 058f66c8a753..6b4ee5b32e0e 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>, tree> *implicit_mappers;
   hash_set<tree> *privatized_types;
   tree clauses;
   /* Iteration variables in an OMP_FOR.  */
@@ -452,6 +453,7 @@ new_omp_context (enum omp_region_type region_type)
   c = XCNEW (struct gimplify_omp_ctx);
   c->outer_context = gimplify_omp_ctxp;
   c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
+  c->implicit_mappers = new hash_map<omp_name_type<tree>, tree>;
   c->privatized_types = new hash_set<tree>;
   c->location = input_location;
   c->region_type = region_type;
@@ -475,6 +477,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
 {
   splay_tree_delete (c->variables);
   delete c->privatized_types;
+  delete c->implicit_mappers;
   c->loop_iter_var.release ();
   XDELETE (c);
 }
@@ -11500,6 +11503,227 @@ error_out:
   return success;
 }
 
+struct instantiate_mapper_info
+{
+  tree *mapper_clauses_p;
+  struct gimplify_omp_ctx *omp_ctx;
+  gimple_seq *pre_p;
+};
+
+/* Helper function for omp_instantiate_mapper.  */
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+  copy_body_data *id = (copy_body_data *) data;
+
+  if (DECL_P (*tp))
+    {
+      tree replacement = remap_decl (*tp, id);
+      if (*tp != replacement)
+	{
+	  *tp = unshare_expr (replacement);
+	  *walk_subtrees = 0;
+	}
+    }
+
+  return NULL_TREE;
+}
+
+/* A copy_decl implementation (for use with tree-inline.cc functions) that
+   only transform decls or SSA names that are part of a map we already
+   prepared.  */
+
+static tree
+omp_mapper_copy_decl (tree var, copy_body_data *cb)
+{
+  tree *repl = cb->decl_map->get (var);
+
+  if (repl)
+    return *repl;
+
+  return var;
+}
+
+static tree *
+omp_instantiate_mapper (gimple_seq *pre_p,
+			hash_map<omp_name_type<tree>, tree> *implicit_mappers,
+			tree mapperfn, tree expr, enum gomp_map_kind outer_kind,
+			tree *mapper_clauses_p)
+{
+  tree mapper_name = NULL_TREE;
+  tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapperfn);
+  gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+  tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+  tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+
+  /* The "extraction map" is used to map the mapper variable in the "declare
+     mapper" directive, and also any temporary variables that have been created
+     as part of expanding the mapper function's body (which are expanded as a
+     "bind" expression in the pre_p sequence).  */
+  hash_map<tree, tree> extraction_map;
+
+  extraction_map.put (dummy_var, expr);
+  extraction_map.put (expr, expr);
+
+  /* This copy_body_data is only used to remap the decls in the
+     OMP_DECLARE_MAPPER tree node expansion itself.  All relevant decls should
+     already be in the current function.  */
+  copy_body_data id;
+  memset (&id, 0, sizeof (id));
+  id.src_fn = current_function_decl;
+  id.dst_fn = current_function_decl;
+  id.src_cfun = cfun;
+  id.decl_map = &extraction_map;
+  id.copy_decl = omp_mapper_copy_decl;
+  id.transform_call_graph_edges = CB_CGE_DUPLICATE; // ???
+  id.transform_new_cfg = true; // ???
+
+  for (; clause; clause = OMP_CLAUSE_CHAIN (clause))
+    {
+      enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (clause);
+      tree *nested_mapper_p = NULL;
+
+      if (map_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+	{
+	  mapper_name = OMP_CLAUSE_DECL (clause);
+	  continue;
+	}
+      else if (map_kind == GOMP_MAP_POP_MAPPER_NAME)
+	{
+	  mapper_name = NULL_TREE;
+	  continue;
+	}
+
+      tree decl = OMP_CLAUSE_DECL (clause);
+      tree unshared, type;
+      bool nonunit_array_with_mapper = false;
+
+      if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+	{
+	  location_t loc = OMP_CLAUSE_LOCATION (clause);
+	  tree tmp = lang_hooks.decls.omp_map_array_section (loc, decl);
+	  if (tmp == decl)
+	    {
+	      unshared = unshare_expr (clause);
+	      nonunit_array_with_mapper = true;
+	      type = TREE_TYPE (TREE_TYPE (decl));
+	    }
+	  else
+	    {
+	      unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+					   OMP_CLAUSE_CODE (clause));
+	      OMP_CLAUSE_DECL (unshared) = tmp;
+	      OMP_CLAUSE_SIZE (unshared)
+		= DECL_P (tmp) ? DECL_SIZE_UNIT (tmp)
+			       : TYPE_SIZE_UNIT (TREE_TYPE (tmp));
+	      type = TREE_TYPE (tmp);
+	    }
+	}
+      else
+	{
+	  unshared = unshare_expr (clause);
+	  type = TREE_TYPE (decl);
+	}
+
+      walk_tree (&unshared, remap_mapper_decl_1, &id, NULL);
+
+      if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
+	OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
+
+      decl = OMP_CLAUSE_DECL (unshared);
+      type = TYPE_MAIN_VARIANT (type);
+
+      nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
+
+      if (nested_mapper_p && *nested_mapper_p != mapperfn)
+	{
+	  if (nonunit_array_with_mapper)
+	    {
+	      sorry ("user-defined mapper with non-unit length array section");
+	      continue;
+	    }
+
+	  if (map_kind == GOMP_MAP_UNSET)
+	    map_kind = outer_kind;
+
+	  mapper_clauses_p
+	    = omp_instantiate_mapper (pre_p, implicit_mappers,
+				      *nested_mapper_p, decl, map_kind,
+				      mapper_clauses_p);
+	  continue;
+	}
+
+      *mapper_clauses_p = unshared;
+      mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared);
+    }
+
+  return mapper_clauses_p;
+}
+
+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;
+  unsigned flags = n->value;
+
+  if (flags & (GOVD_EXPLICIT | GOVD_LOCAL))
+    return 0;
+  if ((flags & GOVD_SEEN) == 0)
+    return 0;
+
+  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 });
+    }
+
+  /* If we already have clauses pertaining to a struct variable, then we don't
+     want to implicitly invoke a user-defined mapper.  */
+  bool handled_p = false;
+  splay_tree_node an
+    = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+  if (an)
+    {
+      unsigned flags = an->value;
+      if ((flags & GOVD_EXPLICIT) != 0
+	  && AGGREGATE_TYPE_P (TREE_TYPE (decl)))
+	handled_p = true;
+    }
+
+  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 (im_info->pre_p, ctx->implicit_mappers,
+				  *mapper_p, decl, GOMP_MAP_TOFROM,
+				  im_info->mapper_clauses_p);
+      /* Make sure we don't map the same variable implicitly in
+	 gimplify_adjust_omp_clauses_1 also.  */
+      n->value |= GOVD_EXPLICIT;
+    }
+
+  return 0;
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
@@ -11511,7 +11735,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   using namespace omp_addr_tokenizer;
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
-  tree *prev_list_p = NULL, *orig_list_p = list_p;
+  tree *orig_list_p = list_p;
   int handled_depend_iterators = -1;
   int nowait = -1;
 
@@ -11542,73 +11766,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	break;
       }
 
-  if (code == OMP_TARGET
-      || code == OMP_TARGET_DATA
-      || code == OMP_TARGET_ENTER_DATA
-      || code == OMP_TARGET_EXIT_DATA)
-    {
-      vec<omp_mapping_group> *groups;
-      groups = omp_gather_mapping_groups (list_p);
-      if (groups)
-	{
-	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
-	  grpmap = omp_index_mapping_groups (groups);
+  vec<omp_mapping_group> *groups = omp_gather_mapping_groups (list_p);
+  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+  unsigned grpnum = 0;
+  tree *grp_start_p = NULL, grp_end = NULL_TREE;
 
-	  omp_resolve_clause_dependencies (code, groups, grpmap);
-	  omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
-					  list_p);
-
-	  omp_mapping_group *outlist = NULL;
-
-	  /* Topological sorting may fail if we have duplicate nodes, which
-	     we should have detected and shown an error for already.  Skip
-	     sorting in that case.  */
-	  if (seen_error ())
-	    goto failure;
-
-	  delete grpmap;
-	  delete groups;
-
-	  /* Rebuild now we have struct sibling lists.  */
-	  groups = omp_gather_mapping_groups (list_p);
-	  grpmap = omp_index_mapping_groups (groups);
-
-	  outlist = omp_tsort_mapping_groups (groups, grpmap);
-	  outlist = omp_segregate_mapping_groups (outlist);
-	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
-
-	failure:
-	  delete grpmap;
-	  delete groups;
-	}
-    }
-  else if (region_type & ORT_ACC)
-    {
-      vec<omp_mapping_group> *groups;
-      groups = omp_gather_mapping_groups (list_p);
-      if (groups)
-	{
-	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
-	  grpmap = omp_index_mapping_groups (groups);
-
-	  oacc_resolve_clause_dependencies (groups, grpmap);
-	  omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
-					  list_p);
-
-	  delete groups;
-	  delete grpmap;
-	}
-    }
+  if (groups)
+    grpmap = omp_index_mapping_groups (groups);
 
   while ((c = *list_p) != NULL)
     {
       bool remove = false;
       bool notice_outer = true;
+      bool map_descriptor;
       const char *check_non_private = NULL;
       unsigned int flags;
       tree decl;
       auto_vec<omp_addr_token *, 10> addr_tokens;
 
+      if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+	{
+	  grp_start_p = NULL;
+	  grp_end = NULL_TREE;
+	}
+
       switch (OMP_CLAUSE_CODE (c))
 	{
 	case OMP_CLAUSE_PRIVATE:
@@ -11913,80 +12094,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  goto do_add;
 
 	case OMP_CLAUSE_MAP:
+	  if (!grp_start_p)
+	    {
+	      grp_start_p = list_p;
+	      grp_end = (*groups)[grpnum].grp_end;
+	      grpnum++;
+	    }
 	  decl = OMP_CLAUSE_DECL (c);
 
+	  if (error_operand_p (decl))
+	    {
+	      remove = true;
+	      break;
+	    }
+
 	  if (!omp_parse_expr (addr_tokens, decl))
 	    {
 	      remove = true;
 	      break;
 	    }
 
-	  if (error_operand_p (decl))
-	    remove = true;
-	  switch (code)
-	    {
-	    case OMP_TARGET:
-	      break;
-	    case OACC_DATA:
-	      if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
-		break;
-	      goto check_firstprivate;
-	    case OACC_ENTER_DATA:
-	    case OACC_EXIT_DATA:
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
-		  && addr_tokens[0]->type == ARRAY_BASE)
-		remove = true;
-	      /* FALLTHRU */
-	    case OMP_TARGET_DATA:
-	    case OMP_TARGET_ENTER_DATA:
-	    case OMP_TARGET_EXIT_DATA:
-	    case OACC_HOST_DATA:
-	    check_firstprivate:
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
-		  || (OMP_CLAUSE_MAP_KIND (c)
-		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
-		/* For target {,enter ,exit }data only the array slice is
-		   mapped, but not the pointer to it.  */
-		remove = true;
-	      break;
-	    default:
-	      break;
-	    }
-	  /* For Fortran, not only the pointer to the data is mapped but also
-	     the address of the pointer, the array descriptor etc.; for
-	     'exit data' - and in particular for 'delete:' - having an 'alloc:'
-	     does not make sense.  Likewise, for 'update' only transferring the
-	     data itself is needed as the rest has been handled in previous
-	     directives.  However, for 'exit data', the array descriptor needs
-	     to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
-
-	     NOTE: Generally, it is not safe to perform "enter data" operations
-	     on arrays where the data *or the descriptor* may go out of scope
-	     before a corresponding "exit data" operation -- and such a
-	     descriptor may be synthesized temporarily, e.g. to pass an
-	     explicit-shape array to a function expecting an assumed-shape
-	     argument.  Performing "enter data" inside the called function
-	     would thus be problematic.  */
 	  if (code == OMP_TARGET_EXIT_DATA
-	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
-	    OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
-					== GOMP_MAP_DELETE
-					? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
-	  else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
-		   && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
-		       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
-	    remove = true;
-	  else if (code == OMP_TARGET_EXIT_DATA
-		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
-		   && OMP_CLAUSE_CHAIN (c)
-		   && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
-		       == OMP_CLAUSE_MAP)
-		   && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			== GOMP_MAP_ATTACH_DETACH)
-		       || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			   == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
-		   && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
-			(OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+	      && OMP_CLAUSE_CHAIN (c)
+	      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
+	      && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		   == GOMP_MAP_ATTACH_DETACH)
+		  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		      == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+	      && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+		   (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
 	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
 
 	  if (remove)
@@ -12007,41 +12144,51 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			      DECL_NAME (decl));
 		}
 	    }
-	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
-	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
-				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
-			     NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+
+	  map_descriptor = false;
+
+	  /* This condition checks if we're mapping an array descriptor that
+	     isn't inside a derived type -- these have special handling, and
+	     are not handled as structs in omp_build_struct_sibling_lists.
+	     See that function for further details.  */
+	  if (*grp_start_p != grp_end
+	      && OMP_CLAUSE_CHAIN (*grp_start_p)
+	      && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
 	    {
-	      remove = true;
-	      break;
-	    }
-	  else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
-		    || (OMP_CLAUSE_MAP_KIND (c)
-			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
-		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
-		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
-	    {
-	      OMP_CLAUSE_SIZE (c)
-		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
-					   false);
-	      if ((region_type & ORT_TARGET) != 0)
-		omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
-				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	      tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+	      if (OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+		  && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET
+		  && DECL_P (OMP_CLAUSE_DECL (grp_mid)))
+		map_descriptor = true;
 	    }
 
-	  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT_UNORD)
-	      && (addr_tokens[0]->type == STRUCTURE_BASE
-		  || addr_tokens[0]->type == ARRAY_BASE)
-	      && addr_tokens[0]->u.structure_base_kind == BASE_DECL)
+	  /* Adding the decl for a struct access: we haven't created
+	     GOMP_MAP_STRUCT nodes yet, so this statement needs to predict
+	     whether they will be created in gimplify_adjust_omp_clauses.  */
+	  if (c == grp_end
+	      && addr_tokens[0]->type == STRUCTURE_BASE
+	      && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+	      && !map_descriptor)
 	    {
 	      gcc_assert (addr_tokens[1]->type == ACCESS_METHOD);
 	      /* If we got to this struct via a chain of pointers, maybe we
 		 want to map it implicitly instead.  */
 	      if (omp_access_chain_p (addr_tokens, 1))
 		break;
+	      omp_mapping_group *wholestruct;
+	      if (!(region_type & ORT_ACC)
+		  && omp_mapped_by_containing_struct (grpmap,
+						      OMP_CLAUSE_DECL (c),
+						      &wholestruct))
+		break;
 	      decl = addr_tokens[1]->expr;
+	      if (splay_tree_lookup (ctx->variables, (splay_tree_key) decl))
+		break;
+	      /* Standalone attach or detach clauses for a struct element
+		 should not inhibit implicit mapping of the whole struct.  */
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		break;
 	      flags = GOVD_MAP | GOVD_EXPLICIT;
 
 	      gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT
@@ -12049,14 +12196,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      goto do_add_decl;
 	    }
 
-	  if (TREE_CODE (decl) == TARGET_EXPR)
-	    {
-	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
-				 is_gimple_lvalue, fb_lvalue)
-		  == GS_ERROR)
-		remove = true;
-	    }
-	  else if (!DECL_P (decl))
+	  if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
 	      if (TREE_CODE (d) == ARRAY_REF)
@@ -12079,56 +12219,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
-	      /* An "attach/detach" operation on an update directive should
-		 behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that
-		 unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
-		 depends on the previous mapping.  */
-	      if (code == OACC_UPDATE
-		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
-		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
 
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	      if (addr_tokens[0]->type == STRUCTURE_BASE
+		  && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+		  && addr_tokens[1]->type == ACCESS_METHOD
+		  && (addr_tokens[1]->u.access_kind == ACCESS_POINTER
+		      || (addr_tokens[1]->u.access_kind
+			  == ACCESS_POINTER_OFFSET))
+		  && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)))
 		{
-		  if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
-		      == ARRAY_TYPE)
-		    remove = true;
-		  else
-		    {
-		      gomp_map_kind k = ((code == OACC_EXIT_DATA
-					  || code == OMP_TARGET_EXIT_DATA)
-					 ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
-		      OMP_CLAUSE_SET_MAP_KIND (c, k);
-		    }
-		}
-
-	      tree cref = decl;
-
-	      while (TREE_CODE (cref) == ARRAY_REF)
-		cref = TREE_OPERAND (cref, 0);
-
-	      if (TREE_CODE (cref) == INDIRECT_REF)
-		cref = TREE_OPERAND (cref, 0);
-
-	      if (TREE_CODE (cref) == COMPONENT_REF)
-		{
-		  tree base = cref;
-		  while (base && !DECL_P (base))
-		    {
-		      tree innerbase = omp_get_base_pointer (base);
-		      if (!innerbase)
-			break;
-		      base = innerbase;
-		    }
-		  if (base
-		      && DECL_P (base)
-		      && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
-		      && POINTER_TYPE_P (TREE_TYPE (base)))
-		    {
-		      splay_tree_node n
-			= splay_tree_lookup (ctx->variables,
-					     (splay_tree_key) base);
-		      n->value |= GOVD_SEEN;
-		    }
+		  tree base = addr_tokens[1]->expr;
+		  splay_tree_node n
+		    = splay_tree_lookup (ctx->variables,
+					 (splay_tree_key) base);
+		  n->value |= GOVD_SEEN;
 		}
 
 	      if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -12239,27 +12343,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		    }
 		}
-	      else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
-				      fb_lvalue) == GS_ERROR)
-		{
-		  remove = true;
-		  break;
-		}
-
-	      if (!remove
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
-		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
-		  && OMP_CLAUSE_CHAIN (c)
-		  && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
-		  && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-		       == GOMP_MAP_ALWAYS_POINTER)
-		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			  == GOMP_MAP_ATTACH_DETACH)
-		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-			  == GOMP_MAP_TO_PSET)))
-		prev_list_p = list_p;
-
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -12267,43 +12350,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
 
-	  if ((code == OMP_TARGET
-	       || code == OMP_TARGET_DATA
-	       || code == OMP_TARGET_ENTER_DATA
-	       || code == OMP_TARGET_EXIT_DATA)
-	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
-	    {
-	      for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
-		   octx = octx->outer_context)
-		{
-		  splay_tree_node n
-		    = splay_tree_lookup (octx->variables,
-					 (splay_tree_key) OMP_CLAUSE_DECL (c));
-		  /* If this is contained in an outer OpenMP region as a
-		     firstprivate value, remove the attach/detach.  */
-		  if (n && (n->value & GOVD_FIRSTPRIVATE))
-		    {
-		      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
-		      goto do_add;
-		    }
-		}
-
-	      enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
-					     ? GOMP_MAP_DETACH
-					     : GOMP_MAP_ATTACH);
-	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
-	    }
-	  else if ((code == OACC_ENTER_DATA
-		    || code == OACC_EXIT_DATA
-		    || code == OACC_PARALLEL)
-		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
-	    {
-	      enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
-					     ? GOMP_MAP_DETACH
-					     : GOMP_MAP_ATTACH);
-	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
-	    }
-
 	  goto do_add;
 
 	case OMP_CLAUSE_AFFINITY:
@@ -12392,6 +12438,17 @@ 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);
+	    ctx->implicit_mappers->put ({ name, type }, fndecl);
+	    remove = true;
+	    break;
+	  }
+
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  flags = GOVD_EXPLICIT;
@@ -12939,6 +12996,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
+  if (groups)
+    {
+      delete grpmap;
+      delete groups;
+    }
+
   ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
 }
@@ -13406,15 +13469,96 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	  }
     }
 
+  if (code == OMP_TARGET
+      || code == OMP_TARGET_DATA
+      || code == OMP_TARGET_ENTER_DATA
+      || code == OMP_TARGET_EXIT_DATA)
+    {
+      tree mapper_clauses = NULL_TREE;
+      instantiate_mapper_info im_info;
+
+      im_info.mapper_clauses_p = &mapper_clauses;
+      im_info.omp_ctx = ctx;
+      im_info.pre_p = pre_p;
+
+      splay_tree_foreach (ctx->variables,
+			  omp_instantiate_implicit_mappers,
+			  (void *) &im_info);
+
+      if (mapper_clauses)
+	{
+	  mapper_clauses
+	    = lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses);
+
+	  /* Stick the implicitly-expanded mapper clauses at the end of the
+	     clause list.  */
+	  tree *tail = list_p;
+	  while (*tail)
+	    tail = &OMP_CLAUSE_CHAIN (*tail);
+	  *tail = mapper_clauses;
+	}
+
+      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);
+
+	  omp_resolve_clause_dependencies (code, groups, grpmap);
+	  omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+					  &grpmap, list_p);
+
+	  omp_mapping_group *outlist = NULL;
+
+	  delete grpmap;
+	  delete groups;
+
+	  /* Rebuild now we have struct sibling lists.  */
+	  groups = omp_gather_mapping_groups (list_p);
+	  grpmap = omp_index_mapping_groups (groups);
+
+	  outlist = omp_tsort_mapping_groups (groups, grpmap);
+	  outlist = omp_segregate_mapping_groups (outlist);
+	  list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+	  delete grpmap;
+	  delete groups;
+	}
+    }
+  else if (ctx->region_type & ORT_ACC)
+    {
+      vec<omp_mapping_group> *groups;
+      groups = omp_gather_mapping_groups (list_p);
+      if (groups)
+	{
+	  hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+	  grpmap = omp_index_mapping_groups (groups);
+
+	  oacc_resolve_clause_dependencies (groups, grpmap);
+	  omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+					  &grpmap, list_p);
+
+	  delete groups;
+	  delete grpmap;
+	}
+    }
+
   tree attach_list = NULL_TREE;
   tree *attach_tail = &attach_list;
 
+  tree *grp_start_p = NULL, grp_end = NULL_TREE;
+
   while ((c = *list_p) != NULL)
     {
       splay_tree_node n;
       bool remove = false;
       bool move_attach = false;
 
+      if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+	grp_end = NULL_TREE;
+
       switch (OMP_CLAUSE_CODE (c))
 	{
 	case OMP_CLAUSE_FIRSTPRIVATE:
@@ -13569,26 +13713,67 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	  break;
 
 	case OMP_CLAUSE_MAP:
-	  if (code == OMP_TARGET_EXIT_DATA
-	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+	  decl = OMP_CLAUSE_DECL (c);
+	  if (!grp_end)
 	    {
+	      grp_start_p = list_p;
+	      grp_end = *omp_group_last (grp_start_p);
+	    }
+	  switch (code)
+	    {
+	    case OMP_TARGET:
+	      break;
+	    case OACC_DATA:
+	      if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+		break;
+	      goto check_firstprivate;
+	    case OACC_ENTER_DATA:
+	    case OACC_EXIT_DATA:
+	    case OMP_TARGET_DATA:
+	    case OMP_TARGET_ENTER_DATA:
+	    case OMP_TARGET_EXIT_DATA:
+	    case OACC_HOST_DATA:
+	    check_firstprivate:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		  || (OMP_CLAUSE_MAP_KIND (c)
+		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		/* For target {,enter ,exit }data only the array slice is
+		   mapped, but not the pointer to it.  */
+		remove = true;
+	      if (code == OMP_TARGET_EXIT_DATA
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+		remove = true;
+	      break;
+	    default:
+	      break;
+	    }
+	  if (remove)
+	    break;
+	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
+				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
+	  gimplify_omp_ctxp = ctx->outer_context;
+	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val,
+			     fb_rvalue) == GS_ERROR)
+	    {
+	      gimplify_omp_ctxp = ctx;
 	      remove = true;
 	      break;
 	    }
-	  /* If we have a target region, we can push all the attaches to the
-	     end of the list (we may have standalone "attach" operations
-	     synthesized for GOMP_MAP_STRUCT nodes that must be processed after
-	     the attachment point AND the pointed-to block have been mapped).
-	     If we have something else, e.g. "enter data", we need to keep
-	     "attach" nodes together with the previous node they attach to so
-	     that separate "exit data" operations work properly (see
-	     libgomp/target.c).  */
-	  if ((ctx->region_type & ORT_TARGET) != 0
-	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-		  || (OMP_CLAUSE_MAP_KIND (c)
-		      == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
-	    move_attach = true;
-	  decl = OMP_CLAUSE_DECL (c);
+	  else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		    || (OMP_CLAUSE_MAP_KIND (c)
+			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+	    {
+	      OMP_CLAUSE_SIZE (c)
+		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
+					   false);
+	      if ((ctx->region_type & ORT_TARGET) != 0)
+		omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	    }
+	  gimplify_omp_ctxp = ctx;
 	  /* Data clauses associated with reductions must be
 	     compatible with present_or_copy.  Warn and adjust the clause
 	     if that is not the case.  */
@@ -13625,7 +13810,54 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	      remove = true;
 	      break;
 	    }
-	  if (!DECL_P (decl))
+	  /* For Fortran, not only the pointer to the data is mapped but also
+	     the address of the pointer, the array descriptor etc.; for
+	     'exit data' - and in particular for 'delete:' - having an 'alloc:'
+	     does not make sense.  Likewise, for 'update' only transferring the
+	     data itself is needed as the rest has been handled in previous
+	     directives.  However, for 'exit data', the array descriptor needs
+	     to be deleted; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
+
+	     NOTE: Generally, it is not safe to perform "enter data" operations
+	     on arrays where the data *or the descriptor* may go out of scope
+	     before a corresponding "exit data" operation -- and such a
+	     descriptor may be synthesized temporarily, e.g. to pass an
+	     explicit-shape array to a function expecting an assumed-shape
+	     argument.  Performing "enter data" inside the called function
+	     would thus be problematic.  */
+	  if (code == OMP_TARGET_EXIT_DATA
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
+	      && grp_start_p
+	      && OMP_CLAUSE_CHAIN (*grp_start_p) == c)
+	    OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*grp_start_p)
+					== GOMP_MAP_DELETE
+					? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
+	  else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
+		   && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
+	    {
+	      remove = true;
+	      break;
+	    }
+	  else if (code == OMP_TARGET_EXIT_DATA
+		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+		   && OMP_CLAUSE_CHAIN (c)
+		   && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+		       == OMP_CLAUSE_MAP)
+		   && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			== GOMP_MAP_ATTACH_DETACH)
+		       || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			   == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+		   && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+			(OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+	  if (TREE_CODE (decl) == TARGET_EXPR)
+	    {
+	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+				 is_gimple_lvalue, fb_lvalue) == GS_ERROR)
+		remove = true;
+	    }
+	  else if (!DECL_P (decl))
 	    {
 	      if ((ctx->region_type & ORT_TARGET) != 0
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
@@ -13648,8 +13880,123 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 			}
 		    }
 		}
+
+	      tree d = decl, *pd;
+	      if (TREE_CODE (d) == ARRAY_REF)
+		{
+		  while (TREE_CODE (d) == ARRAY_REF)
+		    d = TREE_OPERAND (d, 0);
+		  if (TREE_CODE (d) == COMPONENT_REF
+		      && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+		    decl = d;
+		}
+	      pd = &OMP_CLAUSE_DECL (c);
+	      if (d == decl
+		  && TREE_CODE (decl) == INDIRECT_REF
+		  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+		  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+		      == REFERENCE_TYPE)
+		  && (OMP_CLAUSE_MAP_KIND (c)
+		      != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
+		{
+		  pd = &TREE_OPERAND (decl, 0);
+		  decl = TREE_OPERAND (decl, 0);
+		}
+
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+		switch (code)
+		  {
+		  case OACC_ENTER_DATA:
+		  case OACC_EXIT_DATA:
+		    if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
+			== ARRAY_TYPE)
+		      remove = true;
+		    else if (code == OACC_ENTER_DATA)
+		      goto change_to_attach;
+		    /* Fallthrough.  */
+		  case OMP_TARGET_EXIT_DATA:
+		    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH);
+		    break;
+		  case OACC_UPDATE:
+		    /* An "attach/detach" operation on an update directive
+		       should behave as a GOMP_MAP_ALWAYS_POINTER.  Note that
+		       both GOMP_MAP_ATTACH_DETACH and GOMP_MAP_ALWAYS_POINTER
+		       kinds depend on the previous mapping (for non-TARGET
+		       regions).  */
+		    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
+		    break;
+		  default:
+		  change_to_attach:
+		    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH);
+		    if ((ctx->region_type & ORT_TARGET) != 0)
+		      move_attach = true;
+		  }
+	      else if ((ctx->region_type & ORT_TARGET) != 0
+		       && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+			   || (OMP_CLAUSE_MAP_KIND (c)
+			       == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+		move_attach = true;
+
+	      /* If we have e.g. map(struct: *var), don't gimplify the
+		 argument since omp-low.cc wants to see the decl itself.  */
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+		break;
+
+	      /* We've already partly gimplified this in
+		 gimplify_scan_omp_clauses.  Don't do any more.  */
+	      if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
+		break;
+
+	      gimplify_omp_ctxp = ctx->outer_context;
+	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+				 fb_lvalue) == GS_ERROR)
+		remove = true;
+	      gimplify_omp_ctxp = ctx;
+
 	      break;
 	    }
+
+	  if ((code == OMP_TARGET
+	       || code == OMP_TARGET_DATA
+	       || code == OMP_TARGET_ENTER_DATA
+	       || code == OMP_TARGET_EXIT_DATA)
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	    {
+	      bool firstprivatize = false;
+
+	      for (struct gimplify_omp_ctx *octx = ctx->outer_context; octx;
+		   octx = octx->outer_context)
+		{
+		  splay_tree_node n
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) OMP_CLAUSE_DECL (c));
+		  /* If this is contained in an outer OpenMP region as a
+		     firstprivate value, remove the attach/detach.  */
+		  if (n && (n->value & GOVD_FIRSTPRIVATE))
+		    {
+		      firstprivatize = true;
+		      break;
+		    }
+		}
+
+	      enum gomp_map_kind map_kind;
+	      if (firstprivatize)
+		map_kind = GOMP_MAP_FIRSTPRIVATE_POINTER;
+	      else if (code == OMP_TARGET_EXIT_DATA)
+		map_kind = GOMP_MAP_DETACH;
+	      else
+		map_kind = GOMP_MAP_ATTACH;
+	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+	    }
+	  else if ((ctx->region_type & ORT_ACC) != 0
+		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	    {
+	      enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
+					     ? GOMP_MAP_DETACH
+					     : GOMP_MAP_ATTACH);
+	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+	    }
+
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
 	  if ((ctx->region_type & ORT_TARGET) != 0
 	      && !(n->value & GOVD_SEEN)
@@ -13724,6 +14071,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 			  || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
 			      == 0));
 	    }
+
+	  /* If we have a target region, we can push all the attaches to the
+	     end of the list (we may have standalone "attach" operations
+	     synthesized for GOMP_MAP_STRUCT nodes that must be processed after
+	     the attachment point AND the pointed-to block have been mapped).
+	     If we have something else, e.g. "enter data", we need to keep
+	     "attach" nodes together with the previous node they attach to so
+	     that separate "exit data" operations work properly (see
+	     libgomp/target.c).  */
+	  if ((ctx->region_type & ORT_TARGET) != 0
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || (OMP_CLAUSE_MAP_KIND (c)
+		      == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+	    move_attach = true;
+
 	  break;
 
 	case OMP_CLAUSE_TO:
@@ -16930,6 +17292,15 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
   return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
 }
 
+/* Gimplify an OMP_DECLARE_MAPPER node (by just removing it).  */
+
+static enum gimplify_status
+gimplify_omp_declare_mapper (tree *expr_p)
+{
+  *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
@@ -17852,6 +18223,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 0f1bfbf255c9..fa117884fc90 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -85,6 +85,10 @@ extern enum omp_clause_defaultmap_kind lhd_omp_predetermined_mapping (tree);
 extern tree lhd_omp_assignment (tree, tree, tree);
 extern void lhd_omp_finish_clause (tree, gimple_seq *, bool);
 extern tree lhd_omp_array_size (tree, gimple_seq *);
+extern tree lhd_omp_finish_mapper_clauses (tree);
+extern tree lhd_omp_mapper_lookup (tree, tree);
+extern tree lhd_omp_extract_mapper_directive (tree);
+extern tree lhd_omp_map_array_section (location_t, tree);
 struct gimplify_omp_ctx;
 extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
@@ -273,6 +277,11 @@ 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_MAPPER_LOOKUP lhd_omp_mapper_lookup
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
+  lhd_omp_extract_mapper_directive
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION lhd_omp_map_array_section
 #define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false
 #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
 #define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false
@@ -307,6 +316,10 @@ 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_MAPPER_LOOKUP, \
+  LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, \
+  LANG_HOOKS_OMP_MAP_ARRAY_SECTION, \
   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 7d9756759900..649ee4be9429 100644
--- a/gcc/langhooks.cc
+++ b/gcc/langhooks.cc
@@ -642,6 +642,41 @@ lhd_omp_array_size (tree, gimple_seq *)
   return NULL_TREE;
 }
 
+/* Finalize clause list C after expanding custom mappers for implicitly-mapped
+   variables.  */
+
+tree
+lhd_omp_finish_mapper_clauses (tree c)
+{
+  return c;
+}
+
+/* Look up an OpenMP "declare mapper" mapper.  */
+
+tree
+lhd_omp_mapper_lookup (tree, tree)
+{
+  return NULL_TREE;
+}
+
+/* Given the representation used by the front-end to contain a mapper
+   directive, return the statement for the directive itself.  */
+
+tree
+lhd_omp_extract_mapper_directive (tree)
+{
+  return error_mark_node;
+}
+
+/* Return a simplified form for OMP_ARRAY_SECTION argument, or
+   error_mark_node if impossible.  */
+
+tree
+lhd_omp_map_array_section (location_t, tree)
+{
+  return error_mark_node;
+}
+
 /* 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 b1b2b0e10f09..1d5d6aef95a8 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -313,6 +313,22 @@ 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);
+
+  /* Find a mapper in the current parsing context, given a NAME (or
+     NULL_TREE) and TYPE.  */
+  tree (*omp_mapper_lookup) (tree name, tree type);
+
+  /* Return the statement for the mapper directive definition, from the
+     representation used to contain it (e.g. an inline function
+     declaration).  */
+  tree (*omp_extract_mapper_directive) (tree fndecl);
+
+  /* Return a simplified form for OMP_ARRAY_SECTION argument.  */
+  tree (*omp_map_array_section) (location_t, tree t);
+
   /* 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 f9b7ef3426c4..83401f6d9d00 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -152,6 +152,92 @@ get_openacc_privatization_dump_flags ()
 
 extern tree omp_build_component_ref (tree obj, tree field);
 
+template <typename T>
+struct omp_name_type
+{
+  tree name;
+  T type;
+};
+
+template <>
+struct default_hash_traits <omp_name_type<tree> >
+  : typed_noop_remove <omp_name_type<tree> >
+{
+  GTY((skip)) typedef omp_name_type<tree> value_type;
+  GTY((skip)) typedef omp_name_type<tree> compare_type;
+
+  static hashval_t
+  hash (omp_name_type<tree> 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<tree> p)
+  {
+    return p.type == NULL;
+  }
+
+  static bool
+  is_deleted (omp_name_type<tree>)
+  {
+    return false;
+  }
+
+  static bool
+  equal (const omp_name_type<tree> &a, const omp_name_type<tree> &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<tree> &e)
+  {
+    e.type = NULL;
+  }
+};
+
+template <typename T>
+struct omp_mapper_list
+{
+  hash_set<omp_name_type<T>> *seen_types;
+  vec<tree> *mappers;
+
+  omp_mapper_list (hash_set<omp_name_type<T>> *s, vec<tree> *m)
+    : seen_types (s), mappers (m) { }
+
+  void add_mapper (tree name, T type, tree mapperfn)
+  {
+    /* We can't hash a NULL_TREE...  */
+    if (!name)
+      name = void_node;
+
+    omp_name_type<T> n_t = { name, type };
+
+    if (seen_types->contains (n_t))
+      return;
+
+    seen_types->add (n_t);
+    mappers->safe_push (mapperfn);
+  }
+
+  bool contains (tree name, T type)
+  {
+    if (!name)
+      name = void_node;
+
+    return seen_types->contains ({ name, type });
+  }
+};
+
 namespace omp_addr_tokenizer {
 
 /* These are the ways of accessing a variable that have special-case handling
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c
new file mode 100644
index 000000000000..c4d017036c5e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c
@@ -0,0 +1,22 @@
+/* { dg-do compile { target c++ } } */
+
+struct XYZ {
+  int a;
+  int *b;
+  int c;
+};
+
+#pragma omp declare mapper(struct XYZ t)
+/* { dg-error "missing 'map' clause" "" { target c } .-1 } */
+/* { dg-error "missing 'map' clause before end of line" "" { target c++ } .-2 } */
+
+struct ABC {
+  int *a;
+  int b;
+  int c;
+};
+
+#pragma omp declare mapper(struct ABC d) firstprivate(d.b) 
+/* { dg-error "unexpected clause" "" { target c } .-1 } */
+/* { dg-error "expected end of line before '\\(' token" "" { target c } .-2 } */
+/* { dg-error "unexpected clause before '\\(' token" "" { target c++ } .-3 } */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
new file mode 100644
index 000000000000..983d979d68c5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
@@ -0,0 +1,30 @@
+// { dg-do compile { target c++ } }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+#include <stdlib.h>
+
+// Test named mapper invocation.
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+int main (int argc, char *argv[])
+{
+  int N = 1024;
+#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \
+					     map(s.ptr[:N])
+
+  struct S s;
+  s.ptr = (int *) malloc (sizeof (int) * N);
+
+#pragma omp target map(mapper(mapN), tofrom: s)
+// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(alloc: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/c-c++-common/gomp/declare-mapper-4.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c
new file mode 100644
index 000000000000..6d933e4bf6f4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c
@@ -0,0 +1,78 @@
+/* { dg-do compile { target c++ } } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* Check mapper binding clauses.  */
+
+struct Y {
+  int z;
+};
+
+struct Z {
+  int z;
+};
+
+#pragma omp declare mapper (struct Y y) map(tofrom: y)
+#pragma omp declare mapper (struct Z z) map(tofrom: z)
+
+int foo (void)
+{
+  struct Y yy;
+  struct Z zz;
+  int dummy;
+
+#pragma omp target data map(dummy)
+  {
+  #pragma omp target
+    {
+      yy.z++;
+      zz.z++;
+    }
+    yy.z++;
+  }
+  return yy.z;
+}
+
+struct P
+{
+  struct Z *zp;
+};
+
+int bar (void)
+{
+  struct Y yy;
+  struct Z zz;
+  struct P pp;
+  struct Z t;
+  int dummy;
+
+  pp.zp = &t;
+
+#pragma omp declare mapper (struct Y y) map(tofrom: y.z)
+#pragma omp declare mapper (struct 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(struct 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" { target c++ } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" { target c++ } } } */
+
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\)\)} "original" { target c } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\.z\)\)} "original" { target c } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct P,#pragma omp declare mapper \(struct P x\) map\(tofrom:\(x\.zp\)\[0:1\]\) map\(to:x.zp\)\) mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\)} "original" { target c } } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c
new file mode 100644
index 000000000000..f675a8c68902
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c
@@ -0,0 +1,26 @@
+/* { dg-do compile { target c++ } } */
+
+typedef struct S_ {
+  int *myarr;
+  int size;
+} S;
+
+#pragma omp declare mapper (named: struct S_ v) map(to:v.size, v.myarr) \
+						map(tofrom: v.myarr[0:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
+/* { dg-note "'#pragma omp declare mapper \\(named: S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (named: S v) map(to:v.size, v.myarr) \
+					map(tofrom: v.myarr[0:v.size])
+/* { dg-error "redeclaration of 'named' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(named: S\\)'" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (struct S_ v) map(to:v.size, v.myarr) \
+					 map(tofrom: v.myarr[0:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
+/* { dg-note "'#pragma omp declare mapper \\(S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (S v) map(to:v.size, v.myarr) \
+				 map(tofrom: v.myarr[0:v.size])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(S\\)'" "" { target c++ } .-3 } */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c
new file mode 100644
index 000000000000..a2f6c08cdfdd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c
@@ -0,0 +1,23 @@
+/* { dg-do compile { target c++ } } */
+
+int x = 5;
+
+struct Q {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+};
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+
+struct R {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+};
+
+#pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
+/* { dg-error "'y' undeclared" "" { target c } .-1 } */
+/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
+
+int y = 7;
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c
new file mode 100644
index 000000000000..1b1be9dbb666
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c
@@ -0,0 +1,29 @@
+/* { dg-do compile { target c++ } } */
+
+struct Q {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+};
+
+int foo (void)
+{
+  int x = 5;
+  #pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+  return x;
+}
+
+struct R {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+};
+
+int bar (void)
+{
+  #pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
+  /* { dg-error "'y' undeclared" "" { target c } .-1 } */
+  /* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
+  int y = 7;
+  return y;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c
new file mode 100644
index 000000000000..86ddb942072c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c
@@ -0,0 +1,43 @@
+/* { dg-do compile { target c++ } } */
+
+struct Q {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+  int len;
+};
+
+struct R {
+  struct Q qarr[5];
+};
+
+struct R2 {
+  struct Q *qptr;
+};
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr1[0:myq.len]) \
+					  map(myq.arr2[0:myq.len]) \
+					  map(myq.arr3[0:myq.len])
+
+#pragma omp declare mapper (struct R myr) map(myr.qarr[2:3])
+
+#pragma omp declare mapper (struct R2 myr2) map(myr2.qptr[2:3])
+
+int main (int argc, char *argv[])
+{
+  struct R r;
+  struct R2 r2;
+  int N = 256;
+
+#pragma omp target
+/* { dg-message "sorry, unimplemented: user-defined mapper with non-unit length array section" "" { target *-*-* } .-1 } */
+  {
+    for (int i = 2; i < 5; i++)
+      for (int j = 0; j < N; j++)
+	{
+	  r.qarr[i].arr1[j]++;
+	  r2.qptr[i].arr2[j]++;
+	}
+  }
+}
+
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c
new file mode 100644
index 000000000000..54e58426910e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c
@@ -0,0 +1,34 @@
+/* { dg-do compile { target c++ } } */
+
+int x = 5;
+
+struct Q {
+  int *arr1;
+  int *arr2;
+  int *arr3;
+};
+
+int y = 5;
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+/* { dg-note "'#pragma omp declare mapper \\(Q\\)' previously defined here" "" { target c++ } .-2 } */
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:y])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct Q'" "" { target c } .-1 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(Q\\)'" "" { target c++ } .-2 } */
+
+struct R {
+  int *arr1;
+};
+
+void foo (void)
+{
+#pragma omp declare mapper (struct R myr) map(myr.arr1[0:x])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+/* { dg-note "'#pragma omp declare mapper \\(R\\)' previously declared here" "" { target c++ } .-2 } */
+
+#pragma omp declare mapper (struct R myr) map(myr.arr1[0:y])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct R'" "" { target c } .-1 } */
+/* { dg-error "redeclaration of '#pragma omp declare mapper \\(R\\)'" "" { target c++ } .-2 } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index c749db845b0a..77054b6b56aa 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 000000000000..3177d20adbc2
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C
@@ -0,0 +1,58 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+  S s;
+  s.ptr = new int[N];
+  s.size = N;
+
+#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size])
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(default), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(foo), alloc: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(bar), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+  return 0;
+}
+
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } }
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-2.C
new file mode 100644
index 000000000000..7df72c76e2af
--- /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]) // { dg-note "'#pragma omp declare mapper \\(S\\)' previously declared here" }
+
+  /* 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 \\(S\\)'" }
+
+  /* ...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/tree-core.h b/gcc/tree-core.h
index 80b886cc3e4c..d18d26507df9 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -348,6 +348,10 @@ enum omp_clause_code {
   /* OpenMP clause: doacross ({source,sink}:vec).  */
   OMP_CLAUSE_DOACROSS,
 
+  /* 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 8f95f331779e..29da4e521582 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -994,6 +994,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 ();
 	}
@@ -1057,6 +1066,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))
@@ -3833,6 +3859,21 @@ 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, TREE_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);
+      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 81a6ceaf1813..22e884fff601 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -295,6 +295,7 @@ unsigned const char omp_clause_num_ops[] =
   2, /* OMP_CLAUSE_MAP  */
   1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_DOACROSS  */
+  3, /* OMP_CLAUSE__MAPPER_BINDING_  */
   2, /* OMP_CLAUSE__CACHE_  */
   2, /* OMP_CLAUSE_GANG  */
   1, /* OMP_CLAUSE_ASYNC  */
@@ -386,6 +387,7 @@ const char * const omp_clause_code_name[] =
   "map",
   "has_device_addr",
   "doacross",
+  "_mapper_binding_",
   "_cache_",
   "gang",
   "async",
diff --git a/gcc/tree.def b/gcc/tree.def
index f015021e9dcc..7515586822d1 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: Variable decl.
+   Operand 2: List of clauses.
+   The type of the construct is used for the type to be mapped.  */
+DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 3)
+
 /* 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 0d87cec87d38..5d5f77dee3aa 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1532,6 +1532,13 @@ 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_DECL(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1)
+#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \
+  TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2)
+
 #define OMP_SCAN_BODY(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
 #define OMP_SCAN_CLAUSES(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
 
@@ -1972,6 +1979,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 8e16ce5fc33e..77e4520f4acc 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -188,7 +188,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 000000000000..aba4f4265392
--- /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 000000000000..d848fdb73692
--- /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 000000000000..ea9b7ded75b6
--- /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 000000000000..f194e63b5b7f
--- /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 000000000000..0030de8791a0
--- /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 000000000000..14ed10df7025
--- /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 000000000000..ab6320997148
--- /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 000000000000..3818e5264d35
--- /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;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c
new file mode 100644
index 000000000000..b0fa40929fbc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c
@@ -0,0 +1,60 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct {
+  int *arr;
+  int size;
+} B;
+
+#pragma omp declare mapper (mapB : B myb) map(to: myb.size, myb.arr) \
+					  map(tofrom: myb.arr[0:myb.size])
+
+struct A {
+  int *arr1;
+  B *arr2;
+  int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct A var;
+
+  memset (&var, 0, sizeof var);
+  var.arr1 = (int *) calloc (N, sizeof (int));
+  var.arr2 = (B *) malloc (sizeof (B));
+  var.arr2->arr = (int *) calloc (N, sizeof (float));
+  var.arr2->size = N;
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+			  map(tofrom: x.arr1[0:N]) \
+			  map(mapper(mapB), tofrom: x.arr2[0:1])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	{
+	  var.arr1[i]++;
+	  var.arr2->arr[i]++;
+	}
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (var.arr1[i] == 1);
+      assert (var.arr2->arr[i] == 1);
+      assert (var.arr3[i] == 0);
+    }
+
+  free (var.arr1);
+  free (var.arr2->arr);
+  free (var.arr2);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c
new file mode 100644
index 000000000000..b509ddc412c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c
@@ -0,0 +1,59 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct B_tag {
+  int *arr;
+  int size;
+} B;
+
+#pragma omp declare mapper (B myb) map(to: myb.size, myb.arr) \
+				   map(tofrom: myb.arr[0:myb.size])
+
+struct A {
+  int *arr1;
+  B *arr2;
+  int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct A var;
+
+  memset (&var, 0, sizeof var);
+  var.arr1 = (int *) calloc (N, sizeof (int));
+  var.arr2 = (B *) malloc (sizeof (B));
+  var.arr2->arr = (int *) calloc (N, sizeof (int));
+  var.arr2->size = N;
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+			map(tofrom: x.arr1[0:N]) map(tofrom: x.arr2[0:1])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	{
+	  var.arr1[i]++;
+	  var.arr2->arr[i]++;
+	}
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (var.arr1[i] == 1);
+      assert (var.arr2->arr[i] == 1);
+      assert (var.arr3[i] == 0);
+    }
+
+  free (var.arr1);
+  free (var.arr2->arr);
+  free (var.arr2);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c
new file mode 100644
index 000000000000..cf8919c22edf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c
@@ -0,0 +1,87 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct {
+  int *arr;
+  int size;
+} B;
+
+#pragma omp declare mapper (samename : B myb) map(to: myb.size, myb.arr) \
+					      map(tofrom: myb.arr[0:myb.size])
+
+typedef struct {
+  int *arr;
+  int size;
+} C;
+
+
+struct A {
+  int *arr1;
+  B *arr2;
+  C *arr3;
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct A var;
+
+  memset (&var, 0, sizeof var);
+  var.arr1 = (int *) calloc (N, sizeof (int));
+  var.arr2 = (B *) malloc (sizeof (B));
+  var.arr2->arr = (int *) calloc (N, sizeof (int));
+  var.arr2->size = N;
+  var.arr3 = (C *) malloc (sizeof (C));
+  var.arr3->arr = (int *) calloc (N, sizeof (int));
+  var.arr3->size = N;
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+			map(tofrom: x.arr1[0:N]) \
+			map(mapper(samename), tofrom: x.arr2[0:1])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	{
+	  var.arr1[i]++;
+	  var.arr2->arr[i]++;
+	}
+    }
+  }
+
+  {
+    #pragma omp declare mapper (samename : C myc) map(to: myc.size, myc.arr) \
+			map(tofrom: myc.arr[0:myc.size])
+    #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr3) \
+			map(tofrom: x.arr1[0:N]) \
+			map(mapper(samename), tofrom: *x.arr3)
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	{
+	  var.arr1[i]++;
+	  var.arr3->arr[i]++;
+	}
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (var.arr1[i] == 2);
+      assert (var.arr2->arr[i] == 1);
+      assert (var.arr3->arr[i] == 1);
+    }
+
+  free (var.arr1);
+  free (var.arr2->arr);
+  free (var.arr2);
+  free (var.arr3->arr);
+  free (var.arr3);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c
new file mode 100644
index 000000000000..99b7eedad90f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c
@@ -0,0 +1,55 @@
+/* { dg-do run { target c++ } } */
+
+#include <assert.h>
+
+struct T {
+  int a;
+  int b;
+  int c;
+};
+
+void foo (void)
+{
+  struct T x;
+  x.a = x.b = x.c = 0;
+
+#pragma omp target
+  {
+    x.a++;
+    x.c++;
+  }
+
+  assert (x.a == 1);
+  assert (x.b == 0);
+  assert (x.c == 1);
+}
+
+// An identity mapper.  This should do the same thing as the default!
+#pragma omp declare mapper (struct T v) map(v)
+
+void bar (void)
+{
+  struct T x;
+  x.a = x.b = x.c = 0;
+
+#pragma omp target
+  {
+    x.b++;
+  }
+
+#pragma omp target map(x)
+  {
+    x.a++;
+  }
+
+  assert (x.a == 1);
+  assert (x.b == 1);
+  assert (x.c == 0);
+}
+
+int main (int argc, char *argv[])
+{
+  foo ();
+  bar ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c
new file mode 100644
index 000000000000..e7108da25fef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c
@@ -0,0 +1,57 @@
+/* { dg-do run { target c++ } } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+struct Z {
+  int *arr;
+};
+
+void baz (struct Z *zarr, int len)
+{
+#pragma omp declare mapper (struct Z myvar) map(to: myvar.arr) \
+					    map(tofrom: myvar.arr[0:len])
+  zarr[0].arr = (int *) calloc (len, sizeof (int));
+  zarr[5].arr = (int *) calloc (len, sizeof (int));
+
+#pragma omp target map(zarr, *zarr)
+  {
+    for (int i = 0; i < len; i++)
+      zarr[0].arr[i]++;
+  }
+
+#pragma omp target map(zarr, zarr[5])
+  {
+    for (int i = 0; i < len; i++)
+      zarr[5].arr[i]++;
+  }
+
+#pragma omp target map(zarr[5])
+  {
+    for (int i = 0; i < len; i++)
+      zarr[5].arr[i]++;
+  }
+
+#pragma omp target map(zarr, zarr[5:1])
+  {
+    for (int i = 0; i < len; i++)
+      zarr[5].arr[i]++;
+  }
+
+  for (int i = 0; i < len; i++)
+    assert (zarr[0].arr[i] == 1);
+
+  for (int i = 0; i < len; i++)
+    assert (zarr[5].arr[i] == 3);
+
+  free (zarr[5].arr);
+  free (zarr[0].arr);
+}
+
+int
+main (int argc, char *argv[])
+{
+  struct Z myzarr[10];
+  baz (myzarr, 256);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c
new file mode 100644
index 000000000000..9f85df53998a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c
@@ -0,0 +1,62 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+struct A {
+  int *arr1;
+  float *arr2;
+  int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct A var;
+
+  memset (&var, 0, sizeof var);
+  var.arr1 = (int *) calloc (N, sizeof (int));
+  var.arr2 = (float *) calloc (N, sizeof (float));
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr1) \
+					    map(tofrom: x.arr1[0:N])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	var.arr1[i]++;
+    }
+  }
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr2) \
+					    map(tofrom: x.arr2[0:N])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	var.arr2[i]++;
+    }
+  }
+
+  {
+    #pragma omp declare mapper (struct A x) map(tofrom: x.arr3[0:N])
+    #pragma omp target
+    {
+      for (int i = 0; i < N; i++)
+	var.arr3[i]++;
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (var.arr1[i] == 1);
+      assert (var.arr2[i] == 1);
+      assert (var.arr3[i] == 1);
+    }
+
+  free (var.arr1);
+  free (var.arr2);
+}
