diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index a681438cbbe..184fc076388 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,20 @@ 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);
+      OMP_CLAUSE_ITERATORS (c) = iterators;
+    }
 
   parens.skip_until_found_close (parser);
   return nl;
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 30a03f071d8..cca9f1c000c 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15058,7 +15058,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
       /* We've reached the end of a list of expanded nodes.  Reset the group
 	 start pointer.  */
       if (c == grp_sentinel)
-	grp_start_p = NULL;
+	{
+	  if (grp_start_p
+	      && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+	      && OMP_CLAUSE_ITERATORS (*grp_start_p))
+	    for (tree gc = *grp_start_p; gc != grp_sentinel;
+		 gc = OMP_CLAUSE_CHAIN (gc))
+	      OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+	  grp_start_p = NULL;
+	}
 
       switch (OMP_CLAUSE_CODE (c))
 	{
@@ -15805,6 +15813,12 @@ c_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_ITERATORS (c)
+	      && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+	    {
+	      t = error_mark_node;
+	      break;
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -16497,6 +16511,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	pc = &OMP_CLAUSE_CHAIN (c);
     }
 
+  if (grp_start_p
+      && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+      && OMP_CLAUSE_ITERATORS (*grp_start_p))
+    for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+      OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
   if (simdlen
       && safelen
       && tree_int_cst_lt (OMP_CLAUSE_SAFELEN_EXPR (safelen),
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f50534f5f39..79d6e115d16 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41776,16 +41776,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++;
     }
@@ -41793,6 +41811,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);
@@ -41842,10 +41861,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,
@@ -41909,8 +41947,20 @@ 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);
+      OMP_CLAUSE_ITERATORS (c) = iterators;
+    }
 
   return nlist;
 }
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 4f856a9d749..ba5657f7bc2 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7263,7 +7263,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
       /* We've reached the end of a list of expanded nodes.  Reset the group
 	 start pointer.  */
       if (c == grp_sentinel)
-	grp_start_p = NULL;
+	{
+	  if (grp_start_p
+	      && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+	      && OMP_CLAUSE_ITERATORS (*grp_start_p))
+	    for (tree gc = *grp_start_p; gc != grp_sentinel;
+		 gc = OMP_CLAUSE_CHAIN (gc))
+	      OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+	  grp_start_p = NULL;
+	}
 
       switch (OMP_CLAUSE_CODE (c))
 	{
@@ -8484,6 +8492,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_ITERATORS (c)
+	      && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+	    {
+	      t = error_mark_node;
+	      break;
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -9348,6 +9362,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	pc = &OMP_CLAUSE_CHAIN (c);
     }
 
+  if (grp_start_p
+      && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+      && OMP_CLAUSE_ITERATORS (*grp_start_p))
+    for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+      OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
   if (reduction_seen < 0 && (ordered_seen || schedule_seen))
     reduction_seen = -2;
 
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index b5b1f83db8f..6e532d07fcf 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8823,13 +8823,13 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p)
 }
 
 /* Returns a tree expression containing the total iteration count of the
-   iterator clause decl T.  */
+   iterator IT.  */
 
-static tree
-compute_iterator_count (tree t, gimple_seq *pre_p)
+tree
+compute_iterator_count (tree it, gimple_seq *pre_p)
 {
   tree tcnt = size_one_node;
-  for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+  for (; it; it = TREE_CHAIN (it))
     {
       if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL,
 			 is_gimple_val, fb_rvalue) == GS_ERROR
@@ -8899,21 +8899,17 @@ compute_iterator_count (tree t, gimple_seq *pre_p)
    Returns a pointer to the BIND_EXPR_BODY in the innermost loop body.
    LAST_BIND is set to point to the BIND_EXPR containing the whole loop.  */
 
-static tree *
-build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind)
+tree *
+build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind)
 {
-  tree t = OMP_CLAUSE_DECL (c);
-  gcc_assert (OMP_ITERATOR_DECL_P (t));
-
   if (*last_bind)
     gimplify_and_add (*last_bind, pre_p);
-  tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
+  tree block = TREE_VEC_ELT (it, 5);
   *last_bind = build3 (BIND_EXPR, void_type_node,
 		       BLOCK_VARS (block), NULL, block);
   TREE_SIDE_EFFECTS (*last_bind) = 1;
-  SET_EXPR_LOCATION (*last_bind, OMP_CLAUSE_LOCATION (c));
   tree *p = &BIND_EXPR_BODY (*last_bind);
-  for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+  for (; it; it = TREE_CHAIN (it))
     {
       tree var = TREE_VEC_ELT (it, 0);
       tree begin = TREE_VEC_ELT (it, 1);
@@ -9023,7 +9019,7 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
 	  {
 	    if (TREE_PURPOSE (t) != last_iter)
 	      {
-		tree tcnt = compute_iterator_count (t, pre_p);
+		tree tcnt = compute_iterator_count (TREE_PURPOSE (t), pre_p);
 		if (!tcnt)
 		  return 2;
 		last_iter = TREE_PURPOSE (t);
@@ -9181,7 +9177,9 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
 	if (OMP_ITERATOR_DECL_P (t))
 	  {
 	    if (TREE_PURPOSE (t) != last_iter)
-	      last_body = build_iterator_loop (c, pre_p, &last_bind);
+	      last_body = build_iterator_loop (TREE_PURPOSE (t), pre_p,
+					       &last_bind);
+	    SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
 	    last_iter = TREE_PURPOSE (t);
 	    if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR)
 	      {
@@ -12078,6 +12076,51 @@ error_out:
   return success;
 }
 
+/* Callback for walk_tree to find a VAR_DECL (stored in DATA) in the
+   tree TP.  */
+
+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 that the clause C uses all the iterator variables.
+   Return TRUE if there are no errors.  */
+
+static bool
+check_omp_map_iterators (tree c)
+{
+  bool error = false;
+  gcc_assert (OMP_CLAUSE_ITERATORS (c));
+
+  /* Do not check internal map kinds.  */
+  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+    return true;
+
+  for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it))
+    {
+      tree var = TREE_VEC_ELT (it, 0);
+      tree t = walk_tree (&OMP_CLAUSE_DECL (c), 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;
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
@@ -12478,6 +12521,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      break;
 	    }
 
+	  if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c))
+	    {
+	      remove = true;
+	      break;
+	    }
+
 	  if (!omp_parse_expr (addr_tokens, decl))
 	    {
 	      remove = true;
@@ -14168,7 +14217,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_CLAUSE_ITERATORS (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 +14386,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_CLAUSE_ITERATORS (c))
+		break;
+
 	      gimplify_omp_ctxp = ctx->outer_context;
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
 				 fb_lvalue) == GS_ERROR)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index da2051b0279..9cf6e207d1c 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12607,6 +12607,163 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     }
 }
 
+extern tree compute_iterator_count (tree it, gimple_seq *pre_p);
+extern tree *build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind);
+
+struct iterator_loop_info_t
+{
+  tree bind;
+  tree count;
+  tree index;
+  tree *body;
+  tree *iterator;
+  hash_map<tree, tree> elems;
+};
+
+typedef hash_map<tree, iterator_loop_info_t> iterator_loop_map_t;
+
+/* Builds a loop to expand any iterators in clause C, reusing any previously
+   built loops if they use the same set of iterators.  Generated Gimple
+   statements are placed into PRE_P.  Information on the loops is held in
+   LOOPS.  finish_omp_map_iterators must be called before the loops are
+   used. */
+
+static void
+lower_omp_map_iterators (tree c, gimple_seq *pre_p, iterator_loop_map_t *loops)
+{
+  if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+    return;
+
+  bool built_p;
+  iterator_loop_info_t &loop = loops->get_or_insert (OMP_CLAUSE_ITERATORS (c),
+						     &built_p);
+  if (!built_p)
+    {
+      loop.count = compute_iterator_count (OMP_CLAUSE_ITERATORS (c), pre_p);
+      if (!loop.count)
+	return;
+
+      loop.body = build_iterator_loop (OMP_CLAUSE_ITERATORS (c), pre_p,
+				       &loop.bind);
+      loop.index = create_tmp_var (sizetype);
+      SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c));
+      loop.iterator = &OMP_CLAUSE_ITERATORS (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 (loop.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, loop.index,
+			     size_binop (PLUS_EXPR, loop.index, size_int (2)));
+      append_to_statement_list_force (tem, loop.body);
+    }
+
+  /* Create array to hold expanded values.  */
+  tree last_count_2 = size_binop (MULT_EXPR, loop.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");
+    }
+  loop.elems.put (c, elems);
+
+  /* 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, loop.count);
+  gimplify_and_add (tem, pre_p);
+}
+
+/* Set EXPR as the hostaddr expression that should result from the clause C.
+   LOOPS holds the intermediate loop info.  Returns the tree that should be
+   passed as the hostaddr.  */
+
+static tree
+lower_omp_map_iterator_expr (tree expr, tree c, iterator_loop_map_t *loops)
+{
+  if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+    return expr;
+
+  iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c));
+  gcc_assert (loop);
+  tree *elems = loop->elems.get (c);
+
+  /* IN LOOP BODY:  */
+  /* elems[idx] = <expr>;  */
+  tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, loop->index, NULL_TREE,
+		     NULL_TREE);
+  tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, void_type_node,
+			 lhs, expr);
+  append_to_statement_list_force (tem, loop->body);
+
+  return build_fold_addr_expr_with_type (*elems, ptr_type_node);
+}
+
+/* Set SIZE as the size expression that should result from the clause C.
+   LOOPS holds the intermediate loop info.  Returns the tree that should be
+   passed as the clause size.  */
+
+static tree
+lower_omp_map_iterator_size (tree size, tree c, iterator_loop_map_t *loops)
+{
+  if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+    return size;
+
+  iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c));
+  gcc_assert (loop);
+  tree *elems = loop->elems.get (c);
+
+  /* IN LOOP BODY:  */
+  /* elems[idx+1] = size;  */
+  tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems,
+		     size_binop (PLUS_EXPR, loop->index, size_int (1)),
+		     NULL_TREE, NULL_TREE);
+  tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+			 void_type_node, lhs, size);
+  append_to_statement_list_force (tem, loop->body);
+
+  return size_int (SIZE_MAX);
+}
+
+/* Finish building the iterator loops in LOOPS, with generated Gimple
+   statements going in PRE_P.  The loops cannot be amended after this is
+   called.  */
+
+static void
+finish_omp_map_iterators (iterator_loop_map_t *loops, gimple_seq *pre_p)
+{
+  for (iterator_loop_map_t::iterator it = loops->begin ();
+       it != loops->end (); ++it)
+    {
+      iterator_loop_info_t &loop = (*it).second;
+      gimplify_and_add (loop.bind, pre_p);
+
+      for (hash_map<tree, tree>::iterator it2 = loop.elems.begin ();
+	   it2 != loop.elems.end (); ++it2)
+	{
+	  tree clause = (*it2).first;
+	  OMP_CLAUSE_DECL (clause) = (*it2).second;
+	  OMP_CLAUSE_SIZE (clause) = size_int (SIZE_MAX);
+	}
+    }
+}
+
 /* Lower the GIMPLE_OMP_TARGET in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
@@ -12617,7 +12774,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   tree child_fn, t, c;
   gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
   gbind *tgt_bind, *bind, *dep_bind = NULL;
-  gimple_seq tgt_body, olist, ilist, fplist, new_body;
+  gimple_seq tgt_body, olist, iterlist, ilist, fplist, new_body;
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
@@ -12628,6 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   tree deep_map_offset_data = NULL_TREE;
   tree deep_map_offset = NULL_TREE;
 
+  iterator_loop_map_t iterator_loops;
+
   offloaded = is_gimple_omp_offloaded (stmt);
   switch (gimple_omp_target_kind (stmt))
     {
@@ -12706,6 +12865,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   push_gimplify_context ();
   fplist = NULL;
 
+  iterlist = NULL;
   ilist = NULL;
   olist = NULL;
   for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -12761,7 +12921,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    gcc_unreachable ();
 	  }
 #endif
-	  /* FALLTHRU */
+	lower_omp_map_iterators (c, &iterlist, &iterator_loops);
+	/* FALLTHRU */
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
       oacc_firstprivate:
@@ -13190,6 +13351,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			  *p = build_fold_indirect_ref (nd);
 		      }
 		    v = build_fold_addr_expr_with_type (v, ptr_type_node);
+		    v = lower_omp_map_iterator_expr (v, c, &iterator_loops);
 		    gimplify_assign (x, v, &ilist);
 		    nc = NULL_TREE;
 		  }
@@ -13263,12 +13425,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
 		  {
 		    gcc_assert (offloaded);
-		    tree avar
-		      = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
-		    mark_addressable (avar);
-		    gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
-		    talign = DECL_ALIGN_UNIT (avar);
+		    tree avar = build_fold_addr_expr (var);
+		    if (!OMP_CLAUSE_ITERATORS (c))
+		      {
+			tree tmp = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
+			mark_addressable (tmp);
+			gimplify_assign (tmp, avar, &ilist);
+			avar = tmp;
+		      }
+		    talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x)));
 		    avar = build_fold_addr_expr (avar);
+		    avar = lower_omp_map_iterator_expr (avar, c,
+							&iterator_loops);
 		    gimplify_assign (x, avar, &ilist);
 		  }
 		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
@@ -13348,6 +13516,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    if (s == NULL_TREE)
 	      s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
 	    s = fold_convert (size_type_node, s);
+	    s = lower_omp_map_iterator_size (s, c, &iterator_loops);
 	    purpose = size_int (map_idx++);
 	    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
 	    if (TREE_CODE (s) != INTEGER_CST)
@@ -13713,6 +13882,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  DECL_INITIAL (TREE_VEC_ELT (t, 2))
 	    = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
 	}
+      finish_omp_map_iterators (&iterator_loops, &iterlist);
       for (int i = 1; i <= 2; i++)
 	if (deep_map_cnt || !TREE_STATIC (TREE_VEC_ELT (t, i)))
 	  {
@@ -14280,6 +14450,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gimple_omp_set_body (stmt, new_body);
     }
 
+  gsi_insert_seq_before (gsi_p, iterlist, GSI_SAME_STMT);
+
   bind = gimple_build_bind (NULL, NULL,
 			    tgt_bind ? gimple_bind_block (tgt_bind)
 				     : NULL_TREE);
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-map-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c
new file mode 100644
index 00000000000..7d6c8dc6255
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-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-map-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
new file mode 100644
index 00000000000..da14d068f19
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-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-map-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
new file mode 100644
index 00000000000..fb0c761018a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+  #pragma omp target \
+      map(to: x, y) \
+      map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+      map(from: z) \
+      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 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):from:D\.\[0-9\]+" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):to:D\.\[0-9\]+" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):attach:D\.\[0-9\]+" 4 "omplower" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 39e586c808c..be2723dcdae 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -911,6 +911,11 @@ 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,");
+      if (OMP_CLAUSE_ITERATORS (clause))
+	{
+	  dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
+	  pp_colon (pp);
+	}
       switch (OMP_CLAUSE_MAP_KIND (clause))
 	{
 	case GOMP_MAP_ALLOC:
diff --git a/gcc/tree.cc b/gcc/tree.cc
index bc50afca9a3..f12d7b8bb8a 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -266,7 +266,7 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_EXCLUSIVE  */
   2, /* OMP_CLAUSE_FROM  */
   2, /* OMP_CLAUSE_TO  */
-  2, /* OMP_CLAUSE_MAP  */
+  3, /* OMP_CLAUSE_MAP  */
   1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_DOACROSS  */
   2, /* OMP_CLAUSE__CACHE_  */
@@ -11598,6 +11598,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
     case OMP_CLAUSE:
       {
 	int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)];
+	/* Do not walk the iterator operand of OpenMP MAP clauses.  */
+	if (OMP_CLAUSE_HAS_ITERATORS (t))
+	  len--;
 	for (int i = 0; i < len; i++)
 	  WALK_SUBTREE (OMP_CLAUSE_OPERAND (t, i));
 	WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (t));
diff --git a/gcc/tree.h b/gcc/tree.h
index 83075b82cc7..384a5f1f250 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1619,6 +1619,13 @@ class auto_suppress_location_wrappers
   != UNKNOWN_LOCATION)
 #define OMP_CLAUSE_LOCATION(NODE)  (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus
 
+#define OMP_CLAUSE_HAS_ITERATORS(NODE) \
+  (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP)
+#define OMP_CLAUSE_ITERATORS(NODE)					\
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE),	\
+					      OMP_CLAUSE_MAP,		\
+					      OMP_CLAUSE_MAP), 2)
+
 /* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest
    is non-rectangular.  */
 #define OMP_FOR_NON_RECTANGULAR(NODE) \
diff --git a/libgomp/target.c b/libgomp/target.c
index cf62af61f3b..463a162879b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -975,6 +975,105 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
     }
 }
 
+static const char *
+kind_to_name (unsigned short kind)
+{
+  if (GOMP_MAP_IMPLICIT_P (kind))
+    kind &= ~GOMP_MAP_IMPLICIT;
+
+  switch (kind & 0xff)
+    {
+    case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+    case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+    case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+    case GOMP_MAP_TO: return "GOMP_MAP_TO";
+    case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET";
+    case GOMP_MAP_FROM: return "GOMP_MAP_FROM";
+    case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM";
+    case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER";
+    case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+    case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+    default: return "unknown";
+    }
+}
+
+/* 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.  ITERATOR_COUNT holds the iteration count of the
+   iterator that generates each map (0 if not generated from an iterator).
+   HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any
+   merging occurs.  */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+			  void **kinds, size_t **iterator_count)
+{
+  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: %u -> %u\n",
+	      (int) *mapnum, (int) 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;
+  *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+
+  for (size_t i = 0; i < *mapnum; i++)
+    {
+      if ((*sizes)[i] == SIZE_MAX)
+	{
+	  uintptr_t *iterator_array = (*hostaddrs)[i];
+	  size_t count = *iterator_array++;
+	  for (size_t j = 0; j < count; j++)
+	    {
+	      new_hostaddrs[new_idx] = (void *) *iterator_array++;
+	      new_sizes[new_idx] = *iterator_array++;
+	      new_kinds[new_idx] = (*skinds)[i];
+	      (*iterator_count)[new_idx] = j + 1;
+	      gomp_debug (1,
+			  "Expanding map %u <%s>: "
+			  "hostaddrs[%u] = %p, sizes[%u] = %lu\n",
+			  (int) i, kind_to_name (new_kinds[new_idx]),
+			  (int) new_idx, new_hostaddrs[new_idx],
+			  (int) new_idx, (unsigned long) 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];
+	  (*iterator_count)[new_idx] = 0;
+	  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,
@@ -991,6 +1090,11 @@ 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;
+  size_t *iterator_count = NULL;
+  if (short_mapkind)
+    iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+					    &kinds, &iterator_count);
   struct target_mem_desc *tgt
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
   tgt->list_count = mapnum;
@@ -1840,14 +1944,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 
   if (pragma_kind & GOMP_MAP_VARS_TARGET)
     {
+      size_t map_num = 0;
       for (i = 0; i < mapnum; i++)
-	{
-	  cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
-	  gomp_copy_host2dev (devicep, aq,
-			      (void *) (tgt->tgt_start + i * sizeof (void *)),
-			      (void *) &cur_node.tgt_offset, sizeof (void *),
-			      true, cbufp);
-	}
+	if (!iterator_count || iterator_count[i] <= 1)
+	  {
+	    cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
+	    gomp_copy_host2dev (devicep, aq,
+				(void *) (tgt->tgt_start + map_num * sizeof (void *)),
+				(void *) &cur_node.tgt_offset, sizeof (void *),
+				true, cbufp);
+	    map_num++;
+	  }
     }
 
   if (cbufp)
@@ -1879,6 +1986,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
     }
 
   gomp_mutex_unlock (&devicep->lock);
+
+  if (iterators_p)
+    {
+      free (hostaddrs);
+      free (sizes);
+      free (kinds);
+      free (iterator_count);
+    }
+
   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..b3d87f231df
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(to: 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..8569b55ab5b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
@@ -0,0 +1,44 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(alloc: 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..be30fa65d80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
@@ -0,0 +1,56 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(to: 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;
+}
