[v2,3/5] openmp: Add support for iterators in 'target update' clauses (C/C++)

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index f72fca1a711..37c419eb326 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -19305,8 +19305,11 @@  c_parser_omp_clause_device_type (c_parser *parser, tree list)
    to ( variable-list )
 
    OpenMP 5.1:
-   from ( [present :] variable-list )
-   to ( [present :] variable-list ) */
+   from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+   to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+   motion-modifier:
+     present | iterator (iterators-definition)  */
 
 static tree
 c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
@@ -19317,15 +19320,88 @@  c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
   if (!parens.require_open (parser))
     return list;
 
+  int pos = 1, colon_pos = 0;
+  int iterator_length = 0;
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+    {
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+	  == CPP_OPEN_PAREN)
+	{
+	  unsigned int 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;
+	    }
+	}
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos += 2;
+      else
+	pos++;
+      if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
+	{
+	  colon_pos = pos;
+	  break;
+	}
+    }
+
   bool present = false;
-  c_token *token = c_parser_peek_token (parser);
+  tree iterators = NULL_TREE;
 
-  if (token->type == CPP_NAME
-      && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
-      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+  for (pos = 1; pos < colon_pos; pos++)
     {
-      present = true;
-      c_parser_consume_token (parser);
+      c_token *token = c_parser_peek_token (parser);
+
+      if (token->type == CPP_COMMA)
+	{
+	  c_parser_consume_token (parser);
+	  continue;
+	}
+      if (token->type == CPP_NAME)
+	{
+	  const char *name = IDENTIFIER_POINTER (token->value);
+	  if (strcmp (name, "present") == 0)
+	    {
+	      if (present)
+		{
+		  c_parser_error (parser, "too many %<present%> modifiers");
+		  parens.skip_until_found_close (parser);
+		  return list;
+		}
+	      present = true;
+	      c_parser_consume_token (parser);
+	    }
+	  else if (strcmp (name, "iterator") == 0)
+	    {
+	      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;
+	    }
+	  else
+	    {
+	      if (kind == OMP_CLAUSE_TO)
+		c_parser_error (parser, "%<to%> clause with motion modifier "
+				"other than %<iterator%> or %<present%>");
+	      else
+		c_parser_error (parser, "%<from%> clause with motion modifier "
+				"other than %<iterator%> or %<present%>");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	}
+    }
+
+  if (colon_pos)
+    {
+      gcc_assert (pos == colon_pos);
+      gcc_assert (c_parser_next_token_is (parser, CPP_COLON));
       c_parser_consume_token (parser);
     }
 
@@ -19336,6 +19412,19 @@  c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
     for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
       OMP_CLAUSE_MOTION_PRESENT (c) = 1;
 
+  if (iterators)
+    {
+      tree block = pop_scope ();
+      if (iterators == error_mark_node)
+	iterators = NULL_TREE;
+      else
+	TREE_VEC_ELT (iterators, 5) = block;
+    }
+
+  if (iterators)
+    for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
   return nl;
 }
 
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 29947177415..cda0ec6efeb 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41659,8 +41659,11 @@  cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
    to ( variable-list )
 
    OpenMP 5.1:
-   from ( [present :] variable-list )
-   to ( [present :] variable-list ) */
+   from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+   to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+   motion-modifier:
+     present | iterator (iterators-definition)  */
 
 static tree
 cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
@@ -41669,15 +41672,94 @@  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;
 
+  size_t pos = 1, colon_pos = 0;
+  int iterator_length = 0;
+  while (cp_lexer_nth_token_is (parser->lexer, pos, CPP_NAME))
+    {
+      if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN))
+	{
+	  unsigned int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+	  if (n != pos + 1)
+	    {
+	      iterator_length = n - pos;
+	      pos = n - 1;
+	    }
+	}
+      if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_COMMA))
+	pos += 2;
+      else
+	pos++;
+      if (cp_lexer_nth_token_is (parser->lexer, pos, CPP_COLON))
+	{
+	  colon_pos = pos;
+	  break;
+	}
+    }
+
   bool present = false;
-  cp_token *token = cp_lexer_peek_token (parser->lexer);
+  tree iterators = NULL_TREE;
+  for (pos = 1; pos < colon_pos; pos++)
+    {
+      cp_token *token = cp_lexer_peek_token (parser->lexer);
 
-  if (token->type == CPP_NAME
-      && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
-      && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+      if (token->type == CPP_COMMA)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  continue;
+	}
+      if (token->type == CPP_NAME)
+	{
+	  const char *name = IDENTIFIER_POINTER (token->u.value);
+	  if (strcmp (name, "present") == 0)
+	    {
+	      if (present)
+		{
+		  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 = true;
+	      cp_lexer_consume_token (parser->lexer);
+	    }
+	  else if (strcmp (name, "iterator") == 0)
+	    {
+	      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;
+	    }
+	  else
+	    {
+	      if (kind == OMP_CLAUSE_TO)
+		cp_parser_error (parser, "%<to%> clause with motion modifier "
+				 "other than %<iterator%> or %<present%>");
+	      else
+		cp_parser_error (parser, "%<from%> clause with motion modifier "
+				 "other than %<iterator%> or %<present%>");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      return list;
+	    }
+	}
+    }
+
+  if (colon_pos)
     {
-      present = true;
-      cp_lexer_consume_token (parser->lexer);
+      gcc_assert (pos == colon_pos);
+      gcc_assert (cp_lexer_next_token_is (parser->lexer, CPP_COLON));
       cp_lexer_consume_token (parser->lexer);
     }
 
@@ -41686,6 +41768,19 @@  cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
     for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
       OMP_CLAUSE_MOTION_PRESENT (c) = 1;
 
+  if (iterators)
+    {
+      tree block = poplevel (1, 1, 0);
+      if (iterators == error_mark_node)
+	iterators = NULL_TREE;
+      else
+	TREE_VEC_ELT (iterators, 5) = block;
+    }
+
+  if (iterators)
+    for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
   return nl;
 }
 
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 549acf4dfbb..6e938296245 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9387,7 +9387,9 @@  gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
   /* 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
+    if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+	 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
 	&& OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
       all_iterators.add (TREE_PURPOSE (OMP_CLAUSE_DECL (c)));
 
@@ -9397,7 +9399,9 @@  gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
     {
       for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
 	{
-	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	  if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+	       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM)
 	      || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
 	    continue;
 
@@ -13004,6 +13008,10 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE__CACHE_:
 	  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))
 	    {
 	      remove = true;
@@ -13012,17 +13020,19 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
-			     NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+	  if (!last_iterators
+	      && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
+				NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
 	      remove = true;
 	      break;
 	    }
 	  if (!DECL_P (decl))
 	    {
-	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
-				 NULL, is_gimple_lvalue, fb_lvalue)
-		  == GS_ERROR)
+	      if (!last_iterators
+		  && gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
+				    NULL, is_gimple_lvalue, fb_lvalue)
+		     == GS_ERROR)
 		{
 		  remove = true;
 		  break;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
new file mode 100644
index 00000000000..3a64f511da4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
@@ -0,0 +1,20 @@ 
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, float **y)
+{
+  #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2])
+
+  #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2])
+
+  #pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2])
+
+  #pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+  /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */
+
+  #pragma omp target update to (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error ".to. clause with motion modifier other than .iterator. or .present. before .something." } */
+  /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
new file mode 100644
index 00000000000..3789a559b6f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
@@ -0,0 +1,17 @@ 
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+  #pragma omp target update to(iterator(i=0:10): x) /* { dg-error "iterator variable .i. not used in clause expression" }*/
+    ;
+
+  #pragma omp target update from(iterator(i=0:10, j=0:20): x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/
+    ;
+
+  #pragma omp target update to(iterator(i=0:10, j=0:20, k=0:30): 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-update-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
new file mode 100644
index 00000000000..d8672b3a242
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
@@ -0,0 +1,17 @@ 
+/* { 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 update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3], y[i][j][:DIM3])
+  #pragma omp target update from (iterator(i=0:DIM1): 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 "to\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+" 1 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 21ecf94ada5..e43f30818d0 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1099,16 +1099,28 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       pp_string (pp, "from(");
       if (OMP_CLAUSE_MOTION_PRESENT (clause))
 	pp_string (pp, "present:");
-      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
-			 spc, flags, false);
+      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);
+	}
+      dump_generic_node (pp, decl, spc, flags, false);
       goto print_clause_size;
 
     case OMP_CLAUSE_TO:
       pp_string (pp, "to(");
       if (OMP_CLAUSE_MOTION_PRESENT (clause))
 	pp_string (pp, "present:");
-      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
-			 spc, flags, false);
+      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);
+	}
+      dump_generic_node (pp, decl, spc, flags, false);
       goto print_clause_size;
 
     case OMP_CLAUSE__CACHE_:
diff --git a/libgomp/target.c b/libgomp/target.c
index 37a4c539647..c69418f0b78 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2224,6 +2224,7 @@  gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
   size_t i;
   struct splay_tree_key_s cur_node;
   const int typemask = short_mapkind ? 0xff : 0x7;
+  bool iterators_p = false;
 
   if (!devicep)
     return;
@@ -2231,6 +2232,10 @@  gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
   if (mapnum == 0)
     return;
 
+  if (short_mapkind)
+    iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+					    &kinds);
+
   gomp_mutex_lock (&devicep->lock);
   if (devicep->state == GOMP_DEVICE_FINALIZED)
     {
@@ -2324,6 +2329,13 @@  gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 	  }
       }
   gomp_mutex_unlock (&devicep->lock);
+
+  if (iterators_p)
+    {
+      free (hostaddrs);
+      free (sizes);
+      free (kinds);
+    }
 }
 
 static struct gomp_offload_icv_list *
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
new file mode 100644
index 00000000000..5a4cad5c219
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
@@ -0,0 +1,65 @@ 
+/* { dg-do run } */
+
+/* Test target enter data and target update to the 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 sum;
+  int expected = mkarray (x);
+
+  #pragma omp target enter data map(to: x[:DIM1])
+  #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+  #pragma omp target map(from: sum)
+    {
+      sum = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  sum += x[i][j];
+    }
+
+  if (sum != expected)
+    return 1;
+
+  expected = 0;
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      {
+	x[i][j] *= rand ();
+	expected += x[i][j];
+      }
+
+  #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2])
+
+  #pragma omp target map(from: sum)
+    {
+      sum = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  sum += x[i][j];
+    }
+
+  return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
new file mode 100644
index 00000000000..949cc266d84
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
@@ -0,0 +1,57 @@ 
+/* { dg-do run } */
+
+/* Test target enter data and target update from the 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));
+      for (int j = 0; j < DIM2; j++)
+	x[i][j] = 0;
+    }
+}
+
+int main (void)
+{
+  int *x[DIM1];
+  int sum, expected;
+
+  mkarray (x);
+
+  #pragma omp target enter data map(alloc: x[:DIM1])
+  #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+  #pragma omp target 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 + 2);
+	    expected += x[i][j];
+	  }
+    }
+
+  /* Host copy of x should remain unchanged.  */
+  sum = 0;
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      sum += x[i][j];
+  if (sum != 0)
+    return 1;
+
+  #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2])
+
+  /* Host copy should now be updated.  */
+  sum = 0;
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      sum += x[i][j];
+  return sum - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
new file mode 100644
index 00000000000..852635e50f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
@@ -0,0 +1,66 @@ 
+/* { dg-do run } */
+
+/* Test target enter data and target update to the target using map
+   iterators with a function.  */
+
+#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));
+      for (int j = 0; j < DIM2; j++)
+	x[i][j] = rand ();
+    }
+}
+
+int f (int i)
+{
+  return i * 2;
+}
+
+int main (void)
+{
+  int *x[DIM1], x_new[DIM1][DIM2];
+  int sum, expected;
+
+  mkarray (x);
+
+  #pragma omp target enter data map(alloc: x[:DIM1])
+  #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+
+  /* Update x on host.  */
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      {
+	x_new[i][j] = x[i][j];
+	x[i][j] = (i + 1) * (j + 2);
+      }
+
+  /* Update a subset of x on target.  */
+  #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2])
+
+  #pragma omp target map(from: sum)
+    {
+      sum = 0;
+      for (int i = 0; i < DIM1; i++)
+	for (int j = 0; j < DIM2; j++)
+	  sum += x[i][j];
+    }
+
+  /* Calculate expected value on host.  */
+  for (int i = 0; i < DIM1/2; i++)
+    for (int j = 0; j < DIM2; j++)
+      x_new[f (i)][j] = x[f (i)][j];
+
+  expected = 0;
+  for (int i = 0; i < DIM1; i++)
+    for (int j = 0; j < DIM2; j++)
+      expected += x_new[i][j];
+
+  return sum - expected;
+}