diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index aff5af17430..f72fca1a711 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18892,7 +18892,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list)
    map ( [map-type-modifier[,] ...] map-kind: variable-list )
 
    map-type-modifier:
-     always | close */
+     always | close | present | iterator (iterators-definition)  */
 
 static tree
 c_parser_omp_clause_map (c_parser *parser, tree list)
@@ -18907,15 +18907,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 
   int pos = 1;
   int map_kind_pos = 0;
-  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+  int iterator_length = 0;
+  for (;;)
     {
-      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+      c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
+      if (tok->type != CPP_NAME)
+	break;
+
+      const char *p = IDENTIFIER_POINTER (tok->value);
+      c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+      if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN)
+	{
+	  unsigned n = pos + 2;
+	  if (c_parser_check_balanced_raw_token_sequence (parser, &n)
+	      && c_parser_peek_nth_token_raw (parser, n)->type
+		 == CPP_CLOSE_PAREN)
+	    {
+	      iterator_length = n - pos + 1;
+	      pos = n;
+	      next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+	    }
+	}
+
+      if (next_tok->type == CPP_COLON)
 	{
 	  map_kind_pos = pos;
 	  break;
 	}
 
-      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+      if (next_tok->type == CPP_COMMA)
 	pos++;
       pos++;
     }
@@ -18923,6 +18943,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
   int always_modifier = 0;
   int close_modifier = 0;
   int present_modifier = 0;
+  tree iterators = NULL_TREE;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       c_token *tok = c_parser_peek_token (parser);
@@ -18964,10 +18985,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	    }
 	  present_modifier++;
 	}
+      else if (strcmp ("iterator", p) == 0
+	       && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN)
+	{
+	  if (iterators)
+	    {
+	      c_parser_error (parser, "too many %<iterator%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	  iterators = c_parser_omp_iterators (parser);
+	  pos += iterator_length - 1;
+	  continue;
+	}
       else
 	{
 	  c_parser_error (parser, "%<map%> clause with map-type modifier other "
-				  "than %<always%>, %<close%> or %<present%>");
+				  "than %<always%>, %<close%>, %<iterator%> "
+				  "or %<present%>");
 	  parens.skip_until_found_close (parser);
 	  return list;
 	}
@@ -19016,8 +19051,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
 				   true);
 
+  if (iterators)
+    {
+      tree block = pop_scope ();
+      if (iterators == error_mark_node)
+	iterators = NULL_TREE;
+      else
+	TREE_VEC_ELT (iterators, 5) = block;
+    }
+
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+       if (iterators)
+	OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+    }
 
   parens.skip_until_found_close (parser);
   return nl;
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 521c0e85605..d631d95f091 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15811,6 +15811,74 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    auto_vec<omp_addr_token *, 10> addr_tokens;
 
 	    t = OMP_CLAUSE_DECL (c);
+	    if (OMP_ITERATOR_DECL_P (t))
+	      {
+		tree iterators = TREE_PURPOSE (t);
+		if (c_omp_finish_iterators (iterators))
+		  {
+		    t = error_mark_node;
+		    break;
+		  }
+
+		/* Find the end of the group of clauses that use the same
+		   iterator.*/
+		tree end_clause;
+		for (end_clause = c; end_clause;
+		     end_clause = OMP_CLAUSE_CHAIN (end_clause))
+		  {
+		    tree nc = OMP_CLAUSE_CHAIN (end_clause);
+		    /* Remove iterator temporarily.  */
+		    OMP_CLAUSE_DECL (end_clause) =
+		      TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+		    if (!nc
+			|| !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+			|| TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+		      break;
+		  }
+		tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+		/* Temporarily split off the group of clauses with the same
+		   iterator.  */
+		OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+		tree new_clauses = c_finish_omp_clauses (c, ort);
+
+		/* Replace the iterators and splice the new clauses in.  */
+		tree *clause_p = &new_clauses;
+		while (*clause_p)
+		  {
+		    /* Skip unwanted clause types.
+		       FIXME: Is this the right thing to do?  */
+		    bool skip = false;
+		    if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+		      switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+			{
+			case GOMP_MAP_TO:
+			case GOMP_MAP_FROM:
+			case GOMP_MAP_ATTACH:
+			case GOMP_MAP_DETACH:
+			  skip = false;
+			  break;
+			default:
+			  skip = true;
+			  break;
+			}
+		    if (skip)
+		      *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+		    else
+		      {
+			OMP_CLAUSE_DECL (*clause_p)
+			  = build_tree_list (iterators,
+					     OMP_CLAUSE_DECL (*clause_p));
+
+			clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+		      }
+		  }
+		*clause_p = next_clause;
+		*pc = new_clauses;
+		pc = clause_p;
+		continue;
+	      }
+
 	    if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	      {
 		grp_start_p = pc;
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index edfa5a49440..29947177415 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41719,16 +41719,34 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 
   int pos = 1;
   int map_kind_pos = 0;
-  while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
-	 || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
+  int iterator_length = 0;
+  for (;;)
     {
-      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON)
+      cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos);
+      if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE))
+	break;
+
+      cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1);
+      if (tok->type == CPP_NAME
+	  && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
+	  && next_tok->type == CPP_OPEN_PAREN)
+	{
+	  int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+	  if (n != pos + 1)
+	    {
+	      iterator_length = n - pos;
+	      pos = n - 1;
+	      next_tok = cp_lexer_peek_nth_token (parser->lexer, n);
+	    }
+	}
+
+      if (next_tok->type == CPP_COLON)
 	{
 	  map_kind_pos = pos;
 	  break;
 	}
 
-      if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+      if (next_tok->type == CPP_COMMA)
 	pos++;
       pos++;
     }
@@ -41736,6 +41754,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
   bool always_modifier = false;
   bool close_modifier = false;
   bool present_modifier = false;
+  tree iterators = NULL_TREE;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -41785,10 +41804,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 	    }
 	  present_modifier = true;
        }
+      else if (strcmp ("iterator", p) == 0
+	       && cp_lexer_peek_nth_token (parser->lexer, 2)->type
+		  == CPP_OPEN_PAREN)
+	{
+	  if (iterators)
+	    {
+	      cp_parser_error (parser, "too many %<iterator%> modifiers");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	  begin_scope (sk_omp, NULL);
+	  iterators = cp_parser_omp_iterators (parser);
+	  pos += iterator_length - 1;
+	  continue;
+	}
       else
 	{
 	  cp_parser_error (parser, "%<map%> clause with map-type modifier other"
-				   " than %<always%>, %<close%> or %<present%>");
+				   " than %<always%>, %<close%>, %<iterator%>"
+				   " or %<present%>");
 	  cp_parser_skip_to_closing_parenthesis (parser,
 						 /*recovering=*/true,
 						 /*or_comma=*/false,
@@ -41852,8 +41890,21 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
 					  NULL, true);
   finish_scope ();
 
+  if (iterators)
+    {
+      tree block = poplevel (1, 1, 0);
+      if (iterators == error_mark_node)
+	iterators = NULL_TREE;
+      else
+	TREE_VEC_ELT (iterators, 5) = block;
+    }
+
   for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      if (iterators)
+	OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+    }
 
   return nlist;
 }
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 7ecad569900..6222cc2fe87 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8375,6 +8375,71 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    auto_vec<omp_addr_token *, 10> addr_tokens;
 
 	    t = OMP_CLAUSE_DECL (c);
+	    if (OMP_ITERATOR_DECL_P (t))
+	      {
+		tree iterators = TREE_PURPOSE (t);
+		if (cp_omp_finish_iterators (iterators))
+		  {
+		    t = error_mark_node;
+		    break;
+		  }
+
+		/* Find the end of the group of clauses that use the same
+		   iterator.*/
+		tree end_clause;
+		for (end_clause = c; end_clause;
+		     end_clause = OMP_CLAUSE_CHAIN (end_clause))
+		  {
+		    tree nc = OMP_CLAUSE_CHAIN (end_clause);
+		    /* Remove iterator temporarily.  */
+		    OMP_CLAUSE_DECL (end_clause) =
+		      TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+		    if (!nc
+			|| !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+			|| TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+		      break;
+		  }
+		tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+		/* Temporarily split off the group of clauses with the same
+		   iterator.  */
+		OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+		tree new_clauses = finish_omp_clauses (c, ort);
+
+		/* Replace the iterators and splice the new clauses in.  */
+		tree *clause_p = &new_clauses;
+		while (*clause_p)
+		  {
+		    OMP_CLAUSE_DECL (*clause_p)
+		      = build_tree_list (iterators,
+					 OMP_CLAUSE_DECL (*clause_p));
+		    /* Skip unwanted clause types.
+		       FIXME: Is this the right thing to do?  */
+		    bool skip = false;
+		    if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+		      switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+			{
+			case GOMP_MAP_TO:
+			case GOMP_MAP_FROM:
+			case GOMP_MAP_ATTACH:
+			case GOMP_MAP_DETACH:
+			  skip = false;
+			  break;
+			default:
+			  skip = true;
+			  break;
+			}
+		    if (skip)
+		      *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+		    else
+		      clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+		  }
+		*clause_p = next_clause;
+		*pc = new_clauses;
+		pc = clause_p;
+		continue;
+	      }
+
 	    if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	      {
 		grp_start_p = pc;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 8519095adef..549acf4dfbb 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8906,7 +8906,12 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind)
   gcc_assert (OMP_ITERATOR_DECL_P (t));
 
   if (*last_bind)
-    gimplify_and_add (*last_bind, pre_p);
+    {
+      bool saved_into_ssa = gimplify_ctxp->into_ssa;
+      gimplify_ctxp->into_ssa = false;
+      gimplify_and_add (*last_bind, pre_p);
+      gimplify_ctxp->into_ssa = saved_into_ssa;
+    }
   tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
   *last_bind = build3 (BIND_EXPR, void_type_node,
 		       BLOCK_VARS (block), NULL, block);
@@ -9330,6 +9335,166 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
   return 1;
 }
 
+/* Callback for walk_tree to find a VAR_DECL for the given tree.  */
+
+static tree
+find_var_decl (tree *tp, int *, void *data)
+{
+  tree t = *tp;
+
+  if (TREE_CODE (t) == VAR_DECL && t == (tree) data)
+    return t;
+
+  return NULL_TREE;
+}
+
+/* Check for clause decls in iterators that do not use all the iterator
+   variables.  */
+
+static bool
+check_iterator_var_usage (tree c)
+{
+  tree decl = OMP_CLAUSE_DECL (c);
+  bool error = false;
+  gcc_assert (OMP_ITERATOR_DECL_P (decl));
+
+  for (tree it = TREE_PURPOSE (decl); it; it = TREE_CHAIN (it))
+    {
+      tree var = TREE_VEC_ELT (it, 0);
+      tree t = walk_tree (&TREE_VALUE (decl), find_var_decl, var, NULL);
+      if (t == NULL_TREE)
+	t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL);
+      if (t == NULL_TREE)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (c),
+		    "iterator variable %qD not used in clause expression",
+		    var);
+	  error = true;
+	}
+    }
+  return !error;
+}
+
+static void
+gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
+{
+  tree last_iter = NULL_TREE;
+  tree last_bind = NULL_TREE;
+  tree last_count = NULL_TREE;
+  tree last_index = NULL_TREE;
+  tree *last_body = NULL;
+
+  /* Find all map iterators in use.  */
+  hash_set<tree> all_iterators;
+  for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	&& OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+      all_iterators.add (TREE_PURPOSE (OMP_CLAUSE_DECL (c)));
+
+  /* Expand all clauses using the same iterator together.  */
+  for (hash_set<tree>::iterator it = all_iterators.begin ();
+       it != all_iterators.end (); ++it)
+    {
+      for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+	{
+	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	      || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+	    continue;
+
+	  tree t = OMP_CLAUSE_DECL (c);
+	  if (TREE_PURPOSE (t) != *it)
+	    continue;
+
+	  if (!check_iterator_var_usage (c))
+	    continue;
+
+	  if (TREE_PURPOSE (t) != last_iter)
+	    {
+	      tree tcnt = compute_iterator_count (t, pre_p);
+	      if (!tcnt)
+		continue;
+
+	      last_iter = TREE_PURPOSE (t);
+	      last_count = tcnt;
+	      last_body = build_iterator_loop (c, pre_p, &last_bind);
+	      last_index = create_tmp_var (sizetype);
+	      SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
+
+	      /* idx = -1;  */
+	      /* This should be initialized to before the individual elements,
+		 as idx is pre-incremented in the loop body.  */
+	      gimple *g = gimple_build_assign (last_index, size_int (-1));
+	      gimple_seq_add_stmt (pre_p, g);
+
+	      /* IN LOOP BODY: */
+	      /* idx += 2;  */
+	      tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+				     void_type_node, last_index,
+				     size_binop (PLUS_EXPR, last_index,
+						 size_int (2)));
+	      append_to_statement_list_force (tem, last_body);
+	    }
+
+	  /* Create array to hold expanded values.  */
+	  tree last_count_2 = size_binop (MULT_EXPR, last_count, size_int (2));
+	  tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1));
+	  tree elems = NULL_TREE;
+	  if (TREE_CONSTANT (arr_length))
+	    {
+	      tree type = build_array_type (ptr_type_node,
+					    build_index_type (arr_length));
+	      elems = create_tmp_var_raw (type);
+	      TREE_ADDRESSABLE (elems) = 1;
+	      gimple_add_tmp_var (elems);
+	    }
+	  else
+	    {
+	      /* Handle dynamic sizes.  */
+	      sorry ("Dynamic iterator sizes not implemented yet.");
+	    }
+
+	  /* elems[0] = count;  */
+	  tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
+			     NULL_TREE, NULL_TREE);
+	  tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+				 void_type_node, lhs, last_count);
+	  gimplify_and_add (tem, pre_p);
+
+	  /* IN LOOP BODY:  */
+	  /* elems[idx] = &<expr>;  */
+	  lhs = build4 (ARRAY_REF, ptr_type_node, elems, last_index, NULL_TREE,
+			NULL_TREE);
+	  tree rhs = build1 (ADDR_EXPR, ptr_type_node, TREE_VALUE (t));
+	  tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+			    void_type_node, lhs, rhs);
+	  append_to_statement_list_force (tem, last_body);
+
+	  /* elems[idx+1] = OMP_CLAUSE_SIZE (c);  */
+	  lhs = build4 (ARRAY_REF, ptr_type_node, elems,
+			size_binop (PLUS_EXPR, last_index, size_int (1)),
+			NULL_TREE, NULL_TREE);
+	  tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+			    void_type_node, lhs, OMP_CLAUSE_SIZE (c));
+	  append_to_statement_list_force (tem, last_body);
+
+	  /* Replace iterator information.  */
+	  TREE_PURPOSE (t) = make_tree_vec (2);
+	  TREE_VEC_ELT (TREE_PURPOSE (t), 0) = last_iter;
+	  TREE_VEC_ELT (TREE_PURPOSE (t), 1) = elems;
+
+	  OMP_CLAUSE_SIZE (c) = size_int (SIZE_MAX);
+	}
+    }
+
+  if (last_bind)
+    {
+      bool saved_into_ssa = gimplify_ctxp->into_ssa;
+      gimplify_ctxp->into_ssa = false;
+      gimplify_and_add (last_bind, pre_p);
+      gimplify_ctxp->into_ssa = saved_into_ssa;
+    }
+}
+
 /* True if mapping node C maps, or unmaps, a (Fortran) array descriptor.  */
 
 static bool
@@ -9538,6 +9703,22 @@ omp_get_base_pointer (tree expr)
   return NULL_TREE;
 }
 
+/* Return the iterator for a mapping group, or NULL if there isn't one.  */
+
+static tree
+omp_group_iterator (omp_mapping_group *grp)
+{
+  tree c = grp->grp_end;
+  if (!OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+    return NULL_TREE;
+
+  tree iter = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+  if (TREE_VEC_LENGTH (iter) == 2)
+    iter = TREE_VEC_ELT (iter, 0);
+
+  return iter;
+}
+
 /* An attach or detach operation depends directly on the address being
    attached/detached.  Return that address, or none if there are no
    attachments/detachments.  */
@@ -9592,7 +9773,7 @@ omp_get_attachment (omp_mapping_group *grp)
 	  case GOMP_MAP_ATTACH_DETACH:
 	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	  case GOMP_MAP_DETACH:
-	    return OMP_CLAUSE_DECL (node);
+	    return OMP_ITERATOR_CLAUSE_DECL (node);
 
 	  default:
 	    internal_error ("unexpected mapping node");
@@ -9604,7 +9785,7 @@ omp_get_attachment (omp_mapping_group *grp)
       node = OMP_CLAUSE_CHAIN (node);
       if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
 	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
-	return OMP_CLAUSE_DECL (node);
+	return OMP_ITERATOR_CLAUSE_DECL (node);
       else
 	internal_error ("unexpected mapping node");
       return error_mark_node;
@@ -9616,7 +9797,7 @@ omp_get_attachment (omp_mapping_group *grp)
 	return OMP_CLAUSE_DECL (*grp->grp_start);
       if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
 	  || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
-	return OMP_CLAUSE_DECL (*grp->grp_start);
+	return OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
       else
 	internal_error ("unexpected mapping node");
       return error_mark_node;
@@ -9670,7 +9851,9 @@ omp_group_last (tree *start_p)
 		     == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
 		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
-		 || omp_map_clause_descriptor_p (nc)))
+		 || omp_map_clause_descriptor_p (nc))
+	     && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+		== OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc)))
 	{
 	  tree nc2 = OMP_CLAUSE_CHAIN (nc);
 	  if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
@@ -9973,7 +10156,7 @@ omp_index_mapping_groups_1 (hash_map<tree_operand_hash_no_se,
 	   node && j < chained;
 	   node = OMP_CLAUSE_CHAIN (node), j++)
 	{
-	  tree decl = OMP_CLAUSE_DECL (node);
+	  tree decl = OMP_ITERATOR_CLAUSE_DECL (node);
 	  /* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF,
 	     meaning node-hash lookups don't work.  This is a workaround for
 	     that, but ideally we should just create the INDIRECT_REF at
@@ -10160,7 +10343,7 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
 	}
     }
 
-  tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+  tree decl = OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
 
   while (decl)
     {
@@ -10699,7 +10882,7 @@ omp_resolve_clause_dependencies (enum tree_code code,
   FOR_EACH_VEC_ELT (*groups, i, grp)
     {
       tree grp_end = grp->grp_end;
-      tree decl = OMP_CLAUSE_DECL (grp_end);
+      tree decl = OMP_ITERATOR_CLAUSE_DECL (grp_end);
 
       gcc_assert (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP);
 
@@ -10886,7 +11069,9 @@ omp_resolve_clause_dependencies (enum tree_code code,
 	  {
 	    omp_mapping_group *struct_group;
 	    if (omp_mapped_by_containing_struct (grpmap, decl, &struct_group)
-		&& *grp->grp_start == grp_end)
+		&& *grp->grp_start == grp_end
+		&& omp_group_iterator (grp)
+		   == omp_group_iterator (struct_group))
 	      {
 		omp_check_mapping_compatibility (OMP_CLAUSE_LOCATION (grp_end),
 						 struct_group, grp);
@@ -11208,7 +11393,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
   using namespace omp_addr_tokenizer;
   poly_offset_int coffset;
   poly_int64 cbitpos;
-  tree ocd = OMP_CLAUSE_DECL (grp_end);
+  tree ocd = OMP_ITERATOR_CLAUSE_DECL (grp_end);
   bool openmp = !(region_type & ORT_ACC);
   bool target = (region_type & ORT_TARGET) != 0;
   tree *continue_at = NULL;
@@ -11806,7 +11991,7 @@ omp_build_struct_sibling_lists (enum tree_code code,
   FOR_EACH_VEC_ELT (*groups, i, grp)
     {
       tree c = grp->grp_end;
-      tree decl = OMP_CLAUSE_DECL (c);
+      tree decl = OMP_ITERATOR_CLAUSE_DECL (c);
       tree grp_end = grp->grp_end;
       auto_vec<omp_addr_token *> addr_tokens;
       tree sentinel = OMP_CLAUSE_CHAIN (grp_end);
@@ -12091,6 +12276,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   tree c;
   tree *orig_list_p = list_p;
   int handled_depend_iterators = -1;
+  tree last_iterators = NULL_TREE;
   int nowait = -1;
 
   ctx = new_omp_context (region_type);
@@ -12471,6 +12657,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      grpnum++;
 	    }
 	  decl = OMP_CLAUSE_DECL (c);
+	  last_iterators =
+	    OMP_ITERATOR_DECL_P (decl) ? TREE_PURPOSE (decl) : NULL_TREE;
+	  OMP_CLAUSE_DECL (c) = decl = OMP_ITERATOR_CLAUSE_DECL (c);
 
 	  if (error_operand_p (decl))
 	    {
@@ -13388,6 +13577,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 	remove = true;
+      if (last_iterators)
+	OMP_CLAUSE_DECL (c) = build_tree_list (last_iterators,
+					       OMP_CLAUSE_DECL (c));
       if (remove)
 	*list_p = OMP_CLAUSE_CHAIN (c);
       else
@@ -14106,7 +14298,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	  break;
 
 	case OMP_CLAUSE_MAP:
-	  decl = OMP_CLAUSE_DECL (c);
+	  decl = OMP_ITERATOR_CLAUSE_DECL (c);
 	  if (!grp_end)
 	    {
 	      grp_start_p = list_p;
@@ -14168,7 +14360,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 				    : TYPE_SIZE_UNIT (TREE_TYPE (decl));
 	    }
 	  gimplify_omp_ctxp = ctx->outer_context;
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
+	  if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+	    /* Gimplify the OMP_CLAUSE_SIZE later, when the iterator is
+	       gimplified.  */
+	    ;
+	  else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
 				  is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
 	      gimplify_omp_ctxp = ctx;
@@ -14333,6 +14529,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	      if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
 		break;
 
+	      /* Do not gimplify the declaration yet for clauses with
+		 iterators.  */
+	      if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+		break;
+
 	      gimplify_omp_ctxp = ctx->outer_context;
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
 				 fb_lvalue) == GS_ERROR)
@@ -14701,6 +14902,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	  break;
 	}
 
+  gimplify_omp_map_iterators (orig_list_p, pre_p);
+
   gimplify_omp_ctxp = ctx->outer_context;
   delete_omp_context (ctx);
 }
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 4d003f42098..262990ecf9a 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1518,7 +1518,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_MAP:
 	  if (ctx->outer)
 	    scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
-	  decl = OMP_CLAUSE_DECL (c);
+	  decl = OMP_ITERATOR_CLAUSE_DECL (c);
 	  /* If requested, make 'decl' addressable.  */
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c))
@@ -12734,11 +12734,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    gcc_unreachable ();
 	  }
 #endif
+	if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+	    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+	  {
+	    map_cnt++;
+	    continue;
+	  }
 	  /* FALLTHRU */
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
       oacc_firstprivate:
-	var = OMP_CLAUSE_DECL (c);
+	var = OMP_ITERATOR_CLAUSE_DECL (c);
 	if (!DECL_P (var))
 	  {
 	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
@@ -13019,7 +13025,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case OMP_CLAUSE_FROM:
 	  oacc_firstprivate_map:
 	    nc = c;
-	    ovar = OMP_CLAUSE_DECL (c);
+	    ovar = OMP_ITERATOR_CLAUSE_DECL (c);
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		    || (OMP_CLAUSE_MAP_KIND (c)
@@ -13039,6 +13045,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		  {
 		    tree x = build_sender_ref (ovar, ctx);
 		    tree v = ovar;
+		    if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+		      {
+			tree iterator = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+			v = TREE_VEC_ELT (iterator, 1);
+		      }
 		    if (in_reduction_clauses
 			&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			&& OMP_CLAUSE_MAP_IN_REDUCTION (c))
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 014ed35ab41..13e3b58cc92 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,19 +13,19 @@ foo (void)
   #pragma omp target map (to:a)
   ;
 
-  #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
   ;
 
-  #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
   ;
 
-  #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
   ;
 
-  #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
   ;
 
-  #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
   ;
 
 
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c
new file mode 100644
index 00000000000..7d6c8dc6255
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, int **y)
+{
+  #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2])
+    ;
+
+  #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2])
+    ;
+
+  #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* { dg-message "unsupported map expression" } */
+    ;
+
+  #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+    ;
+
+  #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i]) /* { dg-message "unsupported map expression" } */
+    ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c
new file mode 100644
index 00000000000..da14d068f19
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c
@@ -0,0 +1,19 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+  #pragma omp target map(iterator(i=0:10), to: x) /* { dg-error "iterator variable .i. not used in clause expression" } */
+    /* Add a reference to x to ensure that the 'to' clause does not get
+       dropped.  */
+    x[0] = 0;
+
+  #pragma omp target map(iterator(i=0:10, j=0:20), to: x[i]) /* { dg-error "iterator variable .j. not used in clause expression" } */
+    ;
+
+  #pragma omp target map(iterator(i=0:10, j=0:20, k=0:30), to: x[i], y[j], z[k])
+  /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */
+  /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */
+  /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */
+    ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c
new file mode 100644
index 00000000000..22becdda559
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+  #pragma omp target map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+		     map(iterator(i=0:DIM1), from: z[i][:DIM2])
+    ;
+}
+
+/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:from:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:to:" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:attach:" 4 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index ab7ecbfd1ef..21ecf94ada5 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -451,6 +451,19 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags)
   pp_right_paren (pp);
 }
 
+static void
+dump_omp_map_iterators (pretty_printer *pp, tree iter, int spc,
+			dump_flags_t flags)
+{
+  if (TREE_VEC_LENGTH (iter) == 6)
+    dump_omp_iterators (pp, iter, spc, flags);
+  else
+    {
+      dump_omp_iterators (pp, TREE_VEC_ELT (iter, 0), spc, flags);
+      pp_string (pp, ":iterator_array=");
+      dump_generic_node (pp, TREE_VEC_ELT (iter, 1), spc, flags, false);
+    }
+}
 
 /* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN.
 
@@ -461,6 +474,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 {
   const char *name;
   const char *modifier = NULL;
+  tree decl = NULL_TREE;
   switch (OMP_CLAUSE_CODE (clause))
     {
     case OMP_CLAUSE_PRIVATE:
@@ -911,6 +925,13 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       pp_string (pp, "map(");
       if (OMP_CLAUSE_MAP_READONLY (clause))
 	pp_string (pp, "readonly,");
+      decl = OMP_CLAUSE_DECL (clause);
+      if (OMP_ITERATOR_DECL_P (decl))
+	{
+	  dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+	  pp_colon (pp);
+	  decl = TREE_VALUE (decl);
+	}
       switch (OMP_CLAUSE_MAP_KIND (clause))
 	{
 	case GOMP_MAP_ALLOC:
@@ -1025,8 +1046,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	  gcc_unreachable ();
 	}
       pp_colon (pp);
-      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
-			 spc, flags, false);
+      dump_generic_node (pp, decl, spc, flags, false);
      print_clause_size:
       if (OMP_CLAUSE_SIZE (clause))
 	{
diff --git a/gcc/tree.h b/gcc/tree.h
index 83075b82cc7..9fb21a95fbc 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -2162,6 +2162,13 @@ class auto_suppress_location_wrappers
 	 && TREE_PURPOSE (NODE)					\
 	 && TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC)
 
+/* Return the iterator expression if NODE contains an iterator.
+   Return the clause decl if NODE does not.  */
+#define OMP_ITERATOR_CLAUSE_DECL(NODE)				\
+	(OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (NODE))		\
+	 ? TREE_VALUE (OMP_CLAUSE_DECL (NODE))			\
+	 : OMP_CLAUSE_DECL (NODE))
+
 /* In a BLOCK (scope) node:
    Variables declared in the scope NODE.  */
 #define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars)
diff --git a/libgomp/target.c b/libgomp/target.c
index 47ec36928a6..37a4c539647 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -972,6 +972,77 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
     }
 }
 
+/* Map entries containing expanded iterators will be flattened and merged into
+   HOSTADDRS, SIZES and KINDS, and MAPNUM updated.  Returns true if there are
+   any iterators found.  HOSTADDRS, SIZES and KINDS must be freed afterwards
+   if any merging occurs.  */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+			  void **kinds)
+{
+  bool iterator_p = false;
+  size_t map_count = 0;
+  unsigned short **skinds = (unsigned short **) kinds;
+
+  for (size_t i = 0; i < *mapnum; i++)
+    if ((*sizes)[i] == SIZE_MAX)
+      {
+	uintptr_t *iterator_array = (*hostaddrs)[i];
+	map_count += iterator_array[0];
+	iterator_p = true;
+      }
+    else
+      map_count++;
+
+  if (!iterator_p)
+    return false;
+
+  gomp_debug (1,
+	      "Expanding iterator maps - number of map entries: %ld -> %ld\n",
+	      *mapnum, map_count);
+  void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *));
+  size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+  unsigned short *new_kinds
+    = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short));
+  size_t new_idx = 0;
+
+  for (size_t i = 0; i < *mapnum; i++)
+    {
+      if ((*sizes)[i] == SIZE_MAX)
+	{
+	  uintptr_t *iterator_array = (*hostaddrs)[i];
+	  size_t count = iterator_array[0];
+	  for (int j = 1; j < count * 2 + 1; j += 2)
+	    {
+	      new_hostaddrs[new_idx] = (void *) iterator_array[j];
+	      new_sizes[new_idx] = iterator_array[j+1];
+	      new_kinds[new_idx] = (*skinds)[i];
+	      gomp_debug (1,
+			  "Expanding map %ld: "
+			  "hostaddrs[%ld] = %p, sizes[%ld] = %ld\n",
+			  i, new_idx, new_hostaddrs[new_idx],
+			  new_idx, new_sizes[new_idx]);
+	      new_idx++;
+	    }
+	}
+      else
+	{
+	  new_hostaddrs[new_idx] = (*hostaddrs)[i];
+	  new_sizes[new_idx] = (*sizes)[i];
+	  new_kinds[new_idx] = (*skinds)[i];
+	  new_idx++;
+	}
+    }
+
+  *mapnum = map_count;
+  *hostaddrs = new_hostaddrs;
+  *sizes = new_sizes;
+  *kinds = new_kinds;
+
+  return true;
+}
+
 static inline __attribute__((always_inline)) struct target_mem_desc *
 gomp_map_vars_internal (struct gomp_device_descr *devicep,
 			struct goacc_asyncqueue *aq, size_t mapnum,
@@ -988,6 +1059,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
   const int typemask = short_mapkind ? 0xff : 0x7;
   struct splay_tree_s *mem_map = &devicep->mem_map;
   struct splay_tree_key_s cur_node;
+  bool iterators_p = false;
+  if (short_mapkind)
+    iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+					    &kinds);
   struct target_mem_desc *tgt
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
   tgt->list_count = mapnum;
@@ -1876,6 +1951,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
     }
 
   gomp_mutex_unlock (&devicep->lock);
+
+  if (iterators_p)
+    {
+      free (hostaddrs);
+      free (sizes);
+      free (kinds);
+    }
+
   return tgt;
 }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
new file mode 100644
index 00000000000..900a0ba2d64
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
@@ -0,0 +1,44 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+   iterators.  */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[])
+{
+  int expected = 0;
+
+  for (int i = 0; i < DIM1; i++)
+    {
+      x[i] = (int *) malloc (DIM2 * sizeof (int));
+      for (int j = 0; j < DIM2; j++)
+	{
+	  x[i][j] = rand ();
+	  expected += x[i][j];
+	}
+    }
+
+  return expected;
+}
+
+int main (void)
+{
+  int *x[DIM1];
+  int y;
+
+  int expected = mkarray (x);
+
+  #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) map(from: y)
+    {
+      y = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  y += x[i][j];
+    }
+
+  return y - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
new file mode 100644
index 00000000000..bad0f7f17b8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays from target using map
+   iterators.  */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+  for (int i = 0; i < DIM1; i++)
+    x[i] = (int *) malloc (DIM2 * sizeof (int));
+}
+
+int main (void)
+{
+  int *x[DIM1];
+  int y, expected;
+
+  mkarray (x);
+
+  #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \
+		     map(from: expected)
+    {
+      expected = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  {
+	    x[i][j] = (i+1) * (j+1);
+	    expected += x[i][j];
+	  }
+    }
+
+  y = 0;
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      y += x[i][j];
+
+  return y - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
new file mode 100644
index 00000000000..e3da479e6cb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
@@ -0,0 +1,54 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+   iterators, with multiple iterators and function calls in the iterator
+   expression.  */
+
+#include <stdlib.h>
+
+#define DIM1 16
+#define DIM2 15
+
+int mkarrays (int *x[], int *y[])
+{
+  int expected = 0;
+
+  for (int i = 0; i < DIM1; i++)
+    {
+      x[i] = (int *) malloc (DIM2 * sizeof (int));
+      y[i] = (int *) malloc (sizeof (int));
+      *y[i] = rand ();
+      for (int j = 0; j < DIM2; j++)
+	{
+	  x[i][j] = rand ();
+	  expected += x[i][j] * *y[i];
+	}
+    }
+
+  return expected;
+}
+
+int f (int i, int j)
+{
+  return i * 4 + j;
+}
+
+int main (void)
+{
+  int *x[DIM1], *y[DIM1];
+  int sum;
+
+  int expected = mkarrays (x, y);
+
+  #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \
+		     map(iterator(i=0:DIM1), to: y[i][:1]) \
+		     map(from: sum)
+    {
+      sum = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  sum += x[i][j] * y[i][0];
+    }
+
+  return sum - expected;
+}
