diff mbox series

OpenMP: Enable 'declare mapper' mappers for 'target update' directives

Message ID 20230906095504.32204-1-julian@codesourcery.com
State New
Headers show
Series OpenMP: Enable 'declare mapper' mappers for 'target update' directives | expand

Commit Message

Julian Brown Sept. 6, 2023, 9:55 a.m. UTC
This patch enables use of 'declare mapper' for 'target update' directives,
for each of C, C++ and Fortran.

There are some implementation choices here and some
"read-between-the-lines" consequences regarding this functionality,
as follows:

 * It is possible to invoke a mapper which contains clauses that
   don't make sense for a given 'target update' operation.  E.g. if a
   mapper definition specifies a "from:" mapping and the user does "target
   update to(...)" which triggers that mapper, the resulting map kind
   (OpenMP 5.2, "Table 5.3: Map-Type Decay of Map Type Combinations")
   is "alloc" (and for the inverse case "release").  For such cases,
   an unconditional warning is issued and the map clause in question is
   dropped from the mapper expansion.  (Other choices might be to make
   this an error, or to do the same thing but silently, or warn only
   given some special option.)

 * The array-shaping operator is *permitted* for map clauses within
   'declare mapper' definitions.  That is because such mappers may be used
   for 'target update' directives, where the array-shaping operator is
   permitted.  I think that makes sense, depending on the semantic model
   of how and when substitution is supposed to take place, but I couldn't
   find such behaviour explicitly mentioned in the spec (as of 5.2).
   If the mapper is triggered by a different directive ("omp target",
   "omp target data", etc.), an error will be raised.

Support is also added for the "mapper" modifier on to/from clauses for
all three base languages.

This version of the patch incorporates signature changes to OpenMP
variable list parsing functions for C and C++ that are already present
on the og13 branch.  It applies on top of the "infrastructure" support
series:

  https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627895.html

and the lvalue parsing/declare mapper series:

  https://gcc.gnu.org/pipermail/gcc-patches/2023-September/629363.html

and the array-shaping operator and strided/rectangular 'target update'
support series:

  https://gcc.gnu.org/pipermail/gcc-patches/2023-September/629422.html

Tested with offloading to NVPTX. OK?

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

gcc/c-family/
	* c-common.h (c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE
	codes.
	* c-omp.cc (omp_basic_map_kind_name): New function.
	(omp_instantiate_mapper): Add 'target update' support.
	(c_omp_instantiate_mappers): Add 'target update' support.

gcc/c/
	* c-parser.cc (c_parser_omp_variable_list): Add ORT parameter.
	Support array-shaping operator in 'declare mapper' definitions.
	(c_parser_omp_var_list_parens): Add ORT parameter.  Pass to
	c_parser_omp_variable_list.
	(c_parser_oacc_data_clause): Update calls to
	c_parser_omp_var_list_parens.
	(c_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to
	c_parser_omp_variable_list in mapper definitions.
	(c_parser_omp_clause_from_to): Add parsing for mapper modifier.
	(c_parser_omp_target_update): Instantiate mappers.

gcc/cp/
	* parser.cc (cp_parser_omp_var_list_no_open): Add ORT parameter.
	Support array-shaping operator in 'declare mapper' definitions.
	(cp_parser_omp_var_list): Add ORT parameter.
	(cp_parser_oacc_data_clause): Update call to cp_parser_omp_var_list.
	(cp_parser_omp_clause_from_to): Add parsing for mapper modifier.
	(cp_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to
	cp_parser_omp_var_list_no_open in mapper definitions.
	(cp_parser_omp_target_update): Instantiate mappers.

gcc/fortran/
	* openmp.cc (gfc_match_motion_var_list): Add parsing for mapper
	modifier.
	(gfc_match_omp_clauses): Adjust error handling for changes to
	gfc_match_motion_var_list.
	* trans-openmp.cc (gfc_trans_omp_target_update): Instantiate mappers.

gcc/testsuite/
	* c-c++-common/gomp/declare-mapper-17.c: New test.
	* c-c++-common/gomp/declare-mapper-19.c: New test.
	* gfortran.dg/gomp/declare-mapper-24.f90: New test.
	* gfortran.dg/gomp/declare-mapper-26.f90: Uncomment 'target update' part
	of test.
	* gfortran.dg/gomp/declare-mapper-26-p.f90: Likewise.
	* gfortran.dg/gomp/declare-mapper-27.f90: New test.

libgomp/
	* testsuite/libgomp.c-c++-common/declare-mapper-18.c: New test.
	* testsuite/libgomp.fortran/declare-mapper-25.f90: New test.
	* testsuite/libgomp.fortran/declare-mapper-25-p.f90: New test.
	* testsuite/libgomp.fortran/declare-mapper-28.f90: New test.
---
 gcc/c-family/c-common.h                       |   2 +
 gcc/c-family/c-omp.cc                         | 107 ++++++++++-
 gcc/c/c-parser.cc                             | 165 +++++++++++++++--
 gcc/cp/parser.cc                              | 169 ++++++++++++++++--
 gcc/fortran/openmp.cc                         |  86 +++++++--
 gcc/fortran/trans-openmp.cc                   |   7 +-
 .../c-c++-common/gomp/declare-mapper-17.c     |  38 ++++
 .../c-c++-common/gomp/declare-mapper-19.c     |  40 +++++
 .../gfortran.dg/gomp/declare-mapper-24.f90    |  43 +++++
 .../gfortran.dg/gomp/declare-mapper-26-p.f90  |   4 +-
 .../gfortran.dg/gomp/declare-mapper-26.f90    |   4 +-
 .../gfortran.dg/gomp/declare-mapper-27.f90    |  25 +++
 .../libgomp.c-c++-common/declare-mapper-18.c  |  33 ++++
 .../libgomp.fortran/declare-mapper-25-p.f90   |  45 +++++
 .../libgomp.fortran/declare-mapper-25.f90     |  49 +++++
 .../libgomp.fortran/declare-mapper-28.f90     |  38 ++++
 16 files changed, 799 insertions(+), 56 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c
 create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90
 create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c
 create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-25-p.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90
diff mbox series

Patch

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 47de4571d9fc..29723eb0b451 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1276,10 +1276,12 @@  enum c_omp_region_type
   C_ORT_DECLARE_SIMD		= 1 << 2,
   C_ORT_TARGET			= 1 << 3,
   C_ORT_EXIT_DATA		= 1 << 4,
+  C_ORT_UPDATE			= 1 << 5,
   C_ORT_DECLARE_MAPPER		= 1 << 6,
   C_ORT_OMP_DECLARE_SIMD	= C_ORT_OMP | C_ORT_DECLARE_SIMD,
   C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET,
   C_ORT_OMP_EXIT_DATA		= C_ORT_OMP | C_ORT_EXIT_DATA,
+  C_ORT_OMP_UPDATE		= C_ORT_OMP | C_ORT_UPDATE,
   C_ORT_OMP_DECLARE_MAPPER	= C_ORT_OMP | C_ORT_DECLARE_MAPPER,
   C_ORT_ACC_TARGET		= C_ORT_ACC | C_ORT_TARGET
 };
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 0eb9868a33bb..45bf8b5cdcd6 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4374,6 +4374,31 @@  omp_map_decayed_kind (enum gomp_map_kind mapper_kind,
   return omp_join_map_kind (decay_to, force_p, always_p, present_p);
 }
 
+/* Return a name to use for a "basic" map kind, e.g. as output from
+   omp_split_map_kind above.  */
+
+static const char *
+omp_basic_map_kind_name (enum gomp_map_kind kind)
+{
+  switch (kind)
+    {
+    case GOMP_MAP_ALLOC:
+      return "alloc";
+    case GOMP_MAP_TO:
+      return "to";
+    case GOMP_MAP_FROM:
+      return "from";
+    case GOMP_MAP_TOFROM:
+      return "tofrom";
+    case GOMP_MAP_RELEASE:
+      return "release";
+    case GOMP_MAP_DELETE:
+      return "delete";
+    default:
+      gcc_unreachable ();
+    }
+}
+
 /* 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.  */
@@ -4441,7 +4466,9 @@  omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
 
       enum gomp_map_kind decayed_kind
 	= omp_map_decayed_kind (clause_kind, outer_kind,
-				(ort & C_ORT_EXIT_DATA) != 0);
+				(ort & C_ORT_EXIT_DATA) != 0
+				|| (outer_kind == GOMP_MAP_FROM
+				    && (ort & C_ORT_UPDATE) != 0));
       OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind);
 
       type = TYPE_MAIN_VARIANT (type);
@@ -4470,8 +4497,51 @@  omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
 	  continue;
 	}
 
-      *outlist = unshared;
-      outlist = &OMP_CLAUSE_CHAIN (unshared);
+      if (ort & C_ORT_UPDATE)
+	{
+	  bool force_p, always_p, present_p;
+	  decayed_kind
+	    = omp_split_map_kind (decayed_kind, &force_p, &always_p,
+				  &present_p);
+	  /* We don't expect to see these flags here.  */
+	  gcc_assert (!force_p && !always_p);
+	  /* For a "target update" operation, we want to turn the map node
+	     expanded from the mapper back into a OMP_CLAUSE_TO or
+	     OMP_CLAUSE_FROM node.  If we can do neither, emit a warning and
+	     drop the clause.  */
+	  switch (decayed_kind)
+	    {
+	    case GOMP_MAP_TO:
+	    case GOMP_MAP_FROM:
+	      {
+		tree xfer
+		  = build_omp_clause (loc, (decayed_kind == GOMP_MAP_TO
+					    ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM));
+		OMP_CLAUSE_DECL (xfer) = OMP_CLAUSE_DECL (unshared);
+		OMP_CLAUSE_SIZE (xfer) = OMP_CLAUSE_SIZE (unshared);
+		/* For FROM/TO clauses, "present" is represented by a flag.
+		   Set it for the expanded clause here.  */
+		if (present_p)
+		  OMP_CLAUSE_MOTION_PRESENT (xfer) = 1;
+		*outlist = xfer;
+		outlist = &OMP_CLAUSE_CHAIN (xfer);
+	      }
+	      break;
+	    default:
+	      clause_kind
+		= omp_split_map_kind (clause_kind, &force_p, &always_p,
+				      &present_p);
+	      warning_at (loc, 0, "dropping %qs clause during mapper expansion "
+			  "in %<#pragma omp target update%>",
+			  omp_basic_map_kind_name (clause_kind));
+	      inform (OMP_CLAUSE_LOCATION (c), "for map clause here");
+	    }
+	}
+      else
+	{
+	  *outlist = unshared;
+	  outlist = &OMP_CLAUSE_CHAIN (unshared);
+	}
     }
 
   return outlist;
@@ -4489,17 +4559,25 @@  c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
   for (pc = &clauses, c = clauses; c; c = *pc)
     {
       bool using_mapper = false;
+      bool update_p = false, update_present_p = false;
 
       switch (OMP_CLAUSE_CODE (c))
 	{
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  update_p = true;
+	  if (OMP_CLAUSE_MOTION_PRESENT (c))
+	    update_present_p = true;
+	  /* Fallthrough.  */
 	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 (!update_p
+		&& (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);
@@ -4536,9 +4614,22 @@  c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
 		continue;
 	      }
 
-	    enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
-	    if (kind == GOMP_MAP_UNSET)
-	      kind = GOMP_MAP_TOFROM;
+	    enum gomp_map_kind kind;
+	    if (update_p)
+	      {
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO)
+		  kind = update_present_p ? GOMP_MAP_PRESENT_TO
+					  : GOMP_MAP_TO;
+		else
+		  kind = update_present_p ? GOMP_MAP_PRESENT_FROM
+					  : GOMP_MAP_FROM;
+	      }
+	    else
+	      {
+		kind = OMP_CLAUSE_MAP_KIND (c);
+		if (kind == GOMP_MAP_UNSET)
+		  kind = GOMP_MAP_TOFROM;
+	      }
 
 	    type = TYPE_MAIN_VARIANT (type);
 
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 1c28f763c8a0..39bce4974eaf 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -13992,6 +13992,7 @@  static tree
 c_parser_omp_variable_list (c_parser *parser,
 			    location_t clause_loc,
 			    enum omp_clause_code kind, tree list,
+			    enum c_omp_region_type ort = C_ORT_OMP,
 			    bool map_lvalue = false)
 {
   auto_vec<omp_dim> dims;
@@ -14094,7 +14095,9 @@  c_parser_omp_variable_list (c_parser *parser,
 	  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
 	  c_omp_array_section_p = true;
 	  c_omp_array_shaping_op_p
-	    = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM);
+	    = (kind == OMP_CLAUSE_TO
+	       || kind == OMP_CLAUSE_FROM
+	       || ort == C_ORT_OMP_DECLARE_MAPPER);
 	  c_expr expr = c_parser_expr_no_commas (parser, NULL);
 	  if (expr.value != error_mark_node)
 	    mark_exp_read (expr.value);
@@ -14531,7 +14534,8 @@  c_parser_omp_variable_list (c_parser *parser,
 
 static tree
 c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
-			      tree list, bool map_lvalue = false)
+			      tree list, enum c_omp_region_type ort = C_ORT_OMP,
+			      bool map_lvalue = false)
 {
   /* The clauses location.  */
   location_t loc = c_parser_peek_token (parser)->location;
@@ -14539,7 +14543,8 @@  c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
   matching_parens parens;
   if (parens.require_open (parser))
     {
-      list = c_parser_omp_variable_list (parser, loc, kind, list, map_lvalue);
+      list = c_parser_omp_variable_list (parser, loc, kind, list, ort,
+					 map_lvalue);
       parens.skip_until_found_close (parser);
     }
   return list;
@@ -14608,7 +14613,8 @@  c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
       gcc_unreachable ();
     }
   tree nl, c;
-  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, false);
+  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, C_ORT_ACC,
+				     false);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -17762,7 +17768,9 @@  c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
     }
 
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
-				   true);
+				   (kind == GOMP_MAP_UNSET
+				    ? C_ORT_OMP_DECLARE_MAPPER
+				    : C_ORT_OMP), true);
 
   tree last_new = NULL_TREE;
 
@@ -18039,25 +18047,148 @@  c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
   if (!parens.require_open (parser))
     return list;
 
-  bool present = false;
-  c_token *token = c_parser_peek_token (parser);
+  int pos = 1, colon_pos = 0;
 
-  if (token->type == CPP_NAME
-      && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
-      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
     {
-      present = true;
-      c_parser_consume_token (parser);
-      c_parser_consume_token (parser);
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos += 2;
+      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+	       == CPP_OPEN_PAREN)
+	{
+	  unsigned int npos = pos + 2;
+	  if (c_parser_check_balanced_raw_token_sequence (parser, &npos)
+	     && (c_parser_peek_nth_token_raw (parser, npos)->type
+		 == CPP_CLOSE_PAREN))
+	    pos = npos + 1;
+	}
+      else
+	pos++;
+      if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
+	{
+	  colon_pos = pos;
+	  break;
+	}
     }
 
-  tree nl = c_parser_omp_variable_list (parser, loc, kind, list, true);
+  int present_modifier = false;
+  int mapper_modifier = false;
+  tree mapper_name = NULL_TREE;
+
+  for (int pos = 1; pos < colon_pos; ++pos)
+    {
+      c_token *tok = c_parser_peek_token (parser);
+      if (tok->type == CPP_COMMA)
+	{
+	  c_parser_consume_token (parser);
+	  continue;
+	}
+      const char *p = IDENTIFIER_POINTER (tok->value);
+      if (strcmp ("present", p) == 0)
+	{
+	  if (present_modifier)
+	    {
+	      c_parser_error (parser, "too many %<present%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	  present_modifier++;
+	  c_parser_consume_token (parser);
+	}
+      else if (strcmp ("mapper", p) == 0)
+	{
+	  c_parser_consume_token (parser);
+
+	  matching_parens mparens;
+	  if (mparens.require_open (parser))
+	    {
+	      if (mapper_modifier)
+		{
+		  c_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.  */
+		  c_parser_consume_token (parser);
+		  mparens.require_close (parser);
+		  parens.skip_until_found_close (parser);
+		  return list;
+		}
+
+	      tok = c_parser_peek_token (parser);
+
+	      switch (tok->type)
+		{
+		case CPP_NAME:
+		  {
+		    mapper_name = tok->value;
+		    c_parser_consume_token (parser);
+		  }
+		  break;
+
+		case CPP_KEYWORD:
+		  if (tok->keyword == RID_DEFAULT)
+		    {
+		      c_parser_consume_token (parser);
+		      break;
+		    }
+		  /* Fallthrough.  */
+
+		default:
+		  error_at (tok->location,
+			    "expected identifier or %<default%>");
+		  return list;
+		}
+
+	      if (!mparens.require_close (parser))
+		{
+		  parens.skip_until_found_close (parser);
+		  return list;
+		}
+
+	      mapper_modifier++;
+	      pos += 3;
+	    }
+	}
+      else
+	{
+	  c_parser_error (parser, "%<to%> or %<from%> clause with modifier "
+			  "other than %<present%> or %<mapper%>");
+	  parens.skip_until_found_close (parser);
+	  return list;
+	}
+    }
+
+  if (colon_pos)
+    c_parser_require (parser, CPP_COLON, "expected %<:%>");
+
+  tree nl = c_parser_omp_variable_list (parser, loc, kind, list, C_ORT_OMP,
+					true);
   parens.skip_until_found_close (parser);
 
-  if (present)
+  if (present_modifier)
     for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
       OMP_CLAUSE_MOTION_PRESENT (c) = 1;
 
+  if (mapper_name)
+    {
+      tree last_new = NULL_TREE;
+      for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+	last_new = c;
+
+      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) = nl;
+      nl = 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 nl;
 }
 
@@ -22487,7 +22618,9 @@  c_parser_omp_target_update (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
-				"#pragma omp target update");
+				"#pragma omp target update", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_UPDATE);
   bool to_clause = false, from_clause = false;
   for (tree c = clauses;
        c && !to_clause && !from_clause;
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index b961dfec3d32..7b3bda9dd961 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37627,6 +37627,7 @@  struct omp_dim
 static tree
 cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 				tree list, bool *colon,
+				enum c_omp_region_type ort = C_ORT_OMP,
 				bool map_lvalue = false)
 {
   auto_vec<omp_dim> dims;
@@ -37655,7 +37656,8 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	  auto s = make_temp_override (parser->omp_array_section_p, true);
 	  auto o = make_temp_override (parser->omp_array_shaping_op_p,
 				       (kind == OMP_CLAUSE_TO
-					|| kind == OMP_CLAUSE_FROM));
+					|| kind == OMP_CLAUSE_FROM
+					|| ort == C_ORT_OMP_DECLARE_MAPPER));
 	  tree reshaped_to = NULL_TREE;
 	  token = cp_lexer_peek_token (parser->lexer);
 	  location_t loc = token->location;
@@ -38127,10 +38129,11 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 
 static tree
 cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
+			enum c_omp_region_type ort = C_ORT_OMP,
 			bool map_lvalue = false)
 {
   if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
-    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, ort,
 					   map_lvalue);
   return list;
 }
@@ -38198,7 +38201,7 @@  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
       gcc_unreachable ();
     }
   tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
+  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, C_ORT_ACC, false);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -40912,23 +40915,153 @@  cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
 
-  bool present = false;
-  cp_token *token = cp_lexer_peek_token (parser->lexer);
+  int pos = 1;
+  int colon_pos = 0;
 
-  if (token->type == CPP_NAME
-      && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
-      && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+  while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME)
     {
-      present = true;
-      cp_lexer_consume_token (parser->lexer);
-      cp_lexer_consume_token (parser->lexer);
+      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+	pos += 2;
+      else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
+	       == CPP_OPEN_PAREN)
+	pos = cp_parser_skip_balanced_tokens (parser, pos + 1);
+      else
+	pos++;
+      if (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_COLON)
+	{
+	  colon_pos = pos;
+	  break;
+	}
     }
 
-  tree nl = cp_parser_omp_var_list_no_open (parser, kind, list, NULL, true);
-  if (present)
+  bool present_modifier = false;
+  bool mapper_modifier = false;
+  tree mapper_name = NULL_TREE;
+
+  for (int pos = 1; pos < colon_pos; ++pos)
+    {
+      cp_token *tok = cp_lexer_peek_token (parser->lexer);
+      if (tok->type == CPP_COMMA)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  continue;
+	}
+      const char *p = IDENTIFIER_POINTER (tok->u.value);
+      if (strcmp ("present", p) == 0)
+	{
+	  if (present_modifier)
+	    {
+	      cp_parser_error (parser, "too many %<present%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	  present_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, "%<to%> or %<from%> clause with "
+			       "modifier other than %<present%> or %<mapper%>");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	}
+    }
+
+  if (colon_pos)
+    cp_parser_require (parser, CPP_COLON, RT_COLON);
+
+  tree nl = cp_parser_omp_var_list_no_open (parser, kind, list, NULL, C_ORT_OMP,
+					    true);
+  if (present_modifier)
     for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
       OMP_CLAUSE_MOTION_PRESENT (c) = 1;
 
+  if (mapper_name)
+    {
+      tree last_new = NULL_TREE;
+      for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+	last_new = c;
+
+      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) = nl;
+      nl = 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 nl;
 }
 
@@ -41162,7 +41295,9 @@  cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind)
      legally.  */
   begin_scope (sk_omp, NULL);
   nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
-					  NULL, true);
+					  NULL, (kind == GOMP_MAP_UNSET
+						 ? C_ORT_OMP_DECLARE_MAPPER
+						 : C_ORT_OMP), true);
   finish_scope ();
 
   tree last_new = NULL_TREE;
@@ -46120,7 +46255,11 @@  cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
 
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
-				 "#pragma omp target update", pragma_tok);
+				 "#pragma omp target update", pragma_tok,
+				 false);
+  if (!processing_template_decl)
+    clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE);
+  clauses = finish_omp_clauses (clauses, C_ORT_OMP_UPDATE);
   bool to_clause = false, from_clause = false;
   for (tree c = clauses;
        c && !to_clause && !from_clause;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 42d003776ee0..e8b1c7851cfc 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1358,16 +1358,68 @@  gfc_match_motion_var_list (const char *str, gfc_omp_namelist **list,
   if (m != MATCH_YES)
     return m;
 
-  match m_present = gfc_match (" present : ");
+  locus old_loc = gfc_current_locus;
+  int present_modifier = 0;
+  int mapper_modifier = 0;
+  locus second_mapper_locus = old_loc;
+  locus second_present_locus = old_loc;
+  char mapper_id[GFC_MAX_SYMBOL_LEN + 1] = { '\0' };
+
+  for (;;)
+    {
+      locus current_locus = gfc_current_locus;
+      if (gfc_match ("present ") == MATCH_YES)
+	{
+	  if (present_modifier++ == 1)
+	    second_present_locus = current_locus;
+	}
+      else if (gfc_match ("mapper ( ") == MATCH_YES)
+	{
+	  if (mapper_modifier++ == 1)
+	    second_mapper_locus = current_locus;
+	  m = gfc_match (" %n ) ", mapper_id);
+	  if (m != MATCH_YES)
+	    return m;
+	  if (strcmp (mapper_id, "default") == 0)
+	    mapper_id[0] = '\0';
+	}
+      else
+	break;
+      gfc_match (", ");
+    }
+
+  if (gfc_match (" : ") != MATCH_YES)
+    {
+      gfc_current_locus = old_loc;
+      present_modifier = 0;
+      mapper_modifier = 0;
+    }
+
+  if (present_modifier > 1)
+    {
+      gfc_error ("too many %<present%> modifiers at %L", &second_present_locus);
+      return MATCH_ERROR;
+    }
+  if (mapper_modifier > 1)
+    {
+      gfc_error ("too many %<mapper%> modifiers at %L", &second_mapper_locus);
+      return MATCH_ERROR;
+    }
 
   m = gfc_match_omp_variable_list ("", list, false, NULL, headp, true, true);
   if (m != MATCH_YES)
     return m;
-  if (m_present == MATCH_YES)
+  gfc_omp_namelist *n;
+  for (n = **headp; n; n = n->next)
     {
-      gfc_omp_namelist *n;
-      for (n = **headp; n; n = n->next)
+      if (present_modifier)
 	n->u.present_modifier = true;
+
+      if (mapper_id[0] != '\0')
+	{
+	  n->u2.udm = gfc_get_omp_namelist_udm ();
+	  n->u2.udm->mapper_id = gfc_get_string ("%s", mapper_id);
+	}
     }
   return MATCH_YES;
 }
@@ -2705,10 +2757,15 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 					      &c->lists[OMP_LIST_FIRSTPRIVATE],
 					      true) == MATCH_YES)
 	    continue;
-	  if ((mask & OMP_CLAUSE_FROM)
-	      && gfc_match_motion_var_list ("from (", &c->lists[OMP_LIST_FROM],
-					     &head) == MATCH_YES)
-	    continue;
+	  if (mask & OMP_CLAUSE_FROM)
+	    {
+	      m = gfc_match_motion_var_list ("from (", &c->lists[OMP_LIST_FROM],
+					     &head);
+	      if (m == MATCH_YES)
+		continue;
+	      else if (m == MATCH_ERROR)
+		goto error;
+	    }
 	  break;
 	case 'g':
 	  if ((mask & OMP_CLAUSE_GANG)
@@ -3695,10 +3752,15 @@  gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      if (m == MATCH_YES)
 		continue;
 	    }
-	  else if ((mask & OMP_CLAUSE_TO)
-		   && gfc_match_motion_var_list ("to (", &c->lists[OMP_LIST_TO],
-						 &head) == MATCH_YES)
-	    continue;
+	  else if (mask & OMP_CLAUSE_TO)
+	    {
+	      m = gfc_match_motion_var_list ("to (", &c->lists[OMP_LIST_TO],
+						 &head);
+	      if (m == MATCH_YES)
+		continue;
+	      else if (m == MATCH_ERROR)
+		goto error;
+	    }
 	  break;
 	case 'u':
 	  if ((mask & OMP_CLAUSE_UNIFORM)
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 3c4d9ccc5432..365aeb5249a4 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -8770,7 +8770,12 @@  gfc_trans_omp_target_update (gfc_code *code)
   tree stmt, omp_clauses;
 
   gfc_start_block (&block);
-  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+  gfc_omp_clauses *target_update_clauses = code->ext.omp_clauses;
+  gfc_omp_instantiate_mappers (code, target_update_clauses, TOC_OPENMP,
+			       OMP_LIST_TO);
+  gfc_omp_instantiate_mappers (code, target_update_clauses, TOC_OPENMP,
+			       OMP_LIST_FROM);
+  omp_clauses = gfc_trans_omp_clauses (&block, target_update_clauses,
 				       code->loc);
   stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node,
 		     omp_clauses);
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c
new file mode 100644
index 000000000000..ddbb59e4f7fa
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c
@@ -0,0 +1,38 @@ 
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+typedef struct {
+  int a, b, c, d;
+} S;
+
+#pragma omp declare mapper (S s) map(alloc: s.a) map(to: s.b) map(from: s.c) \
+				 map(tofrom: s.d)
+#pragma omp declare mapper (update: S s) map(s.a, s.b, s.c, s.d)
+
+int main()
+{
+  S v;
+#pragma omp target update to(v)
+/* { dg-warning {dropping .from. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */
+/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */
+/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.b\)$} 1 "original" } } */
+#pragma omp target update from(v)
+/* { dg-warning {dropping .to. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */
+/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */
+/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\)$} 1 "original" } } */
+
+#pragma omp target update to(mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.c\) to\(v\.b\) to\(v\.a\)$} 1 "original" } } */
+#pragma omp target update from(mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\) from\(v\.b\) from\(v\.a\)$} 1 "original" } } */
+
+#pragma omp target update to(present, mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update to\(present:v\.d\) to\(present:v\.c\) to\(present:v\.b\) to\(present:v\.a\)$} 2 "original" } } */
+#pragma omp target update from(present, mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update from\(present:v\.d\) from\(present:v\.c\) from\(present:v\.b\) from\(present:v\.a\)$} 2 "original" } } */
+
+#pragma omp target update to(present: v.a, v.b, v.c, v.d)
+#pragma omp target update from(present: v.a, v.b, v.c, v.d)
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c
new file mode 100644
index 000000000000..fd40c6a25e8d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c
@@ -0,0 +1,40 @@ 
+/* { dg-do compile } */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <assert.h>
+
+typedef struct {
+  int *ptr;
+} S;
+
+int main(void)
+{
+#pragma omp declare mapper(grid: S x) map(([9][11]) x.ptr[3:3:2][1:4:3])
+  S q;
+  q.ptr = (int *) calloc (9 * 11, sizeof (int));
+
+  /* The 'grid' mapper specifies a noncontiguous region, so it can't be used
+     for 'map' like this.  */
+#pragma omp target enter data map(mapper(grid), to: q)
+/* { dg-error {array section is not contiguous in .map. clause} "" { target *-*-* } .-1 } */
+/* { dg-error {.#pragma omp target enter data. must contain at least one .map. clause} "" { target *-*-* } .-2 } */
+
+#pragma omp target
+  for (int i = 0; i < 9*11; i++)
+    q.ptr[i] = i;
+
+  /* It's OK on a 'target update' directive though.  */
+#pragma omp target update from(mapper(grid): q)
+
+  for (int j = 0; j < 9; j++)
+    for (int i = 0; i < 11; i++)
+      if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0
+	  && i >= 1 && i <= 10 && ((i - 1) % 3) == 0)
+	assert (q.ptr[j * 11 + i] == j * 11 + i);
+
+#pragma omp target exit data map(mapper(grid), release: q)
+/* { dg-error {array section is not contiguous in .map. clause} "" { target *-*-* } .-1 } */
+/* { dg-error {.#pragma omp target exit data. must contain at least one .map. clause} "" { target *-*-* } .-2 } */
+  return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90
new file mode 100644
index 000000000000..9555a94badab
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90
@@ -0,0 +1,43 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+type t
+integer :: a, b, c, d
+end type t
+
+type(t) :: tvar
+
+!$omp declare mapper (T :: t) map(alloc: t%a) map(to: t%b) map(from: t%c) &
+!$omp & map(tofrom: t%d)
+
+!$omp declare mapper (updatey: T :: t) map(t%a) map(t%b) map(t%c) map(t%d)
+
+!$omp target update to(tvar)
+! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 }
+! { dg-warning "Dropping incompatible .FROM. mapper clause" "" { target *-*-* } .-2 }
+! { dg-final { scan-tree-dump-times {(?n)update to\(tvar\.b \[len: [0-9]+\]\) to\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+!$omp target update from(tvar)
+! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 }
+! { dg-warning "Dropping incompatible .TO. mapper clause" "" { target *-*-* } .-2 }
+! { dg-final { scan-tree-dump-times {(?n)update from\(tvar\.c \[len: [0-9]+\]\) from\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+
+!$omp target update to(present: tvar)
+! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 }
+! { dg-warning "Dropping incompatible .FROM. mapper clause" "" { target *-*-* } .-2 }
+! { dg-final { scan-tree-dump-times {(?n)update to\(present:tvar\.b \[len: [0-9]+\]\) to\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+!$omp target update from(present: tvar)
+! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 }
+! { dg-warning "Dropping incompatible .TO. mapper clause" "" { target *-*-* } .-2 }
+! { dg-final { scan-tree-dump-times {(?n)update from\(present:tvar\.c \[len: [0-9]+\]\) from\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+
+!$omp target update to(mapper(updatey): tvar)
+! { dg-final { scan-tree-dump-times {(?n)update to\(tvar\.a \[len: [0-9]+\]\) to\(tvar\.b \[len: [0-9]+\]\) to\(tvar\.c \[len: [0-9]+\]\) to\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+!$omp target update from(mapper(updatey): tvar)
+! { dg-final { scan-tree-dump-times {(?n)update from\(tvar\.a \[len: [0-9]+\]\) from\(tvar\.b \[len: [0-9]+\]\) from\(tvar\.c \[len: [0-9]+\]\) from\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+
+!$omp target update to(present, mapper(updatey): tvar)
+! { dg-final { scan-tree-dump-times {(?n)update to\(present:tvar\.a \[len: [0-9]+\]\) to\(present:tvar\.b \[len: [0-9]+\]\) to\(present:tvar\.c \[len: [0-9]+\]\) to\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+!$omp target update from(present, mapper(updatey): tvar)
+! { dg-final { scan-tree-dump-times {(?n)update from\(present:tvar\.a \[len: [0-9]+\]\) from\(present:tvar\.b \[len: [0-9]+\]\) from\(present:tvar\.c \[len: [0-9]+\]\) from\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26-p.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26-p.f90
index 705decea0fa3..a1b2cf66d572 100644
--- a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26-p.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26-p.f90
@@ -19,8 +19,8 @@  var%arr = 0
 
 var%arr = 1
 
-! But this is fine.  (Re-enabled by a later patch.)
-!!$omp target update to(mapper(even): var)
+! But this is fine.
+!$omp target update to(mapper(even): var)
 
 ! As 'enter data'.
 !$omp target exit data map(mapper(even), delete: var)
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90
index be5605ce8b7b..e81518cfa553 100644
--- a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90
@@ -23,8 +23,8 @@  var%arr = 0
 
 var%arr = 1
 
-! But this is fine.  (Re-enabled by later patch.)
-!!$omp target update to(mapper(even): var)
+! But this is fine.
+!$omp target update to(mapper(even): var)
 
 ! As 'enter data'.
 !$omp target exit data map(mapper(even), delete: var)
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90
new file mode 100644
index 000000000000..6b3a181acaa5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90
@@ -0,0 +1,25 @@ 
+! { dg-do compile }
+
+type t
+integer :: x
+end type t
+
+type(t) :: var
+
+! Error on attempt to use missing named mapper.
+!$omp target update to(mapper(boo): var)
+! { dg-error {User-defined mapper .boo. not found} "" { target *-*-* } .-1 }
+
+var%x = 0
+
+!$omp target map(mapper(boo), tofrom: var)
+! { dg-error {User-defined mapper .boo. not found} "" { target *-*-* } .-1 }
+var%x = 5
+!$omp end target
+
+! These should be fine though...
+!$omp target enter data map(mapper(default), to: var)
+
+!$omp target exit data map(from: var)
+
+end
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c
new file mode 100644
index 000000000000..50f37cba89db
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c
@@ -0,0 +1,33 @@ 
+#include <stdlib.h>
+#include <stdio.h>
+#include <assert.h>
+
+typedef struct {
+  int *ptr;
+} S;
+
+int main(void)
+{
+#pragma omp declare mapper(grid: S x) map(([9][11]) x.ptr[3:3:2][1:4:3])
+  S q;
+  q.ptr = (int *) calloc (9 * 11, sizeof (int));
+
+#pragma omp target enter data map(to: q.ptr, q.ptr[0:9*11])
+
+#pragma omp target
+  for (int i = 0; i < 9*11; i++)
+    q.ptr[i] = i;
+
+#pragma omp target update from(mapper(grid): q)
+
+  for (int j = 0; j < 9; j++)
+    for (int i = 0; i < 11; i++)
+      if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0
+	  && i >= 1 && i <= 10 && ((i - 1) % 3) == 0)
+	assert (q.ptr[j * 11 + i] == j * 11 + i);
+      else
+	assert (q.ptr[j * 11 + i] == 0);
+
+#pragma omp target exit data map(release: q.ptr, q.ptr[0:9*11])
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-25-p.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-25-p.f90
new file mode 100644
index 000000000000..04605b91e475
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-25-p.f90
@@ -0,0 +1,45 @@ 
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+type t
+integer, pointer :: arr(:)
+end type t
+
+!$omp declare mapper(odd: T :: tv) map(tv%arr(1::2))
+!$omp declare mapper(even: T :: tv) map(tv%arr(2::2))
+
+type(t) :: var
+integer :: i
+integer, target :: tgtarr(100)
+
+var%arr => tgtarr
+
+var%arr = 0
+
+!$omp target enter data map(to: var%arr, var%arr(1:100))
+
+var%arr = 1
+
+!$omp target update to(mapper(odd): var)
+
+!$omp target
+do i=1,100
+  if (mod(i,2).eq.0.and.var%arr(i).ne.0) stop 1
+  if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 2
+end do
+!$omp end target
+
+var%arr = 2
+
+!$omp target update to(mapper(even): var)
+
+!$omp target
+do i=1,100
+  if (mod(i,2).eq.0.and.var%arr(i).ne.2) stop 3
+  if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 4
+end do
+!$omp end target
+
+!$omp target exit data map(delete: var%arr, var%arr(1:100))
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90
new file mode 100644
index 000000000000..9c9c805bfddf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90
@@ -0,0 +1,49 @@ 
+! NOTE: Make a 'run' test after allocatable component mappings are fixed.
+! { dg-do compile }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+type t
+integer, allocatable :: arr(:)
+end type t
+
+!$omp declare mapper(odd: T :: tv) map(tv%arr(1::2))
+! { dg-error "List item 'tv' with allocatable components is not permitted in map clause" "" { target *-*-* } .-1 }
+!$omp declare mapper(even: T :: tv) map(tv%arr(2::2))
+! { dg-error "List item 'tv' with allocatable components is not permitted in map clause" "" { target *-*-* } .-1 }
+
+type(t) :: var
+integer :: i
+
+allocate(var%arr(100))
+
+var%arr = 0
+
+!$omp target enter data map(to: var)
+! { dg-error "List item 'var' with allocatable components is not permitted in map clause" "" { target *-*-* } .-1 }
+
+var%arr = 1
+
+!$omp target update to(mapper(odd): var)
+
+!$omp target
+do i=1,100
+  if (mod(i,2).eq.0.and.var%arr(i).ne.0) stop 1
+  if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 2
+end do
+!$omp end target
+
+var%arr = 2
+
+!$omp target update to(mapper(even): var)
+
+!$omp target
+do i=1,100
+  if (mod(i,2).eq.0.and.var%arr(i).ne.2) stop 3
+  if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 4
+end do
+!$omp end target
+
+!$omp target exit data map(delete: var)
+! { dg-error "List item 'var' with allocatable components is not permitted in map clause" "" { target *-*-* } .-1 }
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90
new file mode 100644
index 000000000000..6561decc49a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90
@@ -0,0 +1,38 @@ 
+! { dg-do run }
+
+program p
+
+type t
+integer :: x, y
+end type t
+
+type(t) :: var
+
+var%x = 0
+var%y = 0
+
+var = sub(7)
+
+contains
+
+type(t) function sub(arg)
+integer :: arg
+
+!$omp declare mapper (t :: tvar) map(tvar%x, tvar%y)
+
+!$omp target enter data map(alloc: sub)
+
+sub%x = 5
+sub%y = arg
+
+!$omp target update to(sub)
+
+!$omp target
+if (sub%x.ne.5) stop 1
+if (sub%y.ne.7) stop 2
+!$omp end target
+
+!$omp target exit data map(release: sub)
+
+end function sub
+end program p