openmp: Add support for non-constant iterator parameters in map, to and from clauses

This patch enables support for using non-constant expressions when specifying
iterators in the map clause of target constructs and to/from clauses of
target update constructs.

gcc/

	* gimplify.cc (omp_iterator_elems_length): New.
	(build_omp_iterators_loops): Change type of elements
	array to pointer of pointers if array length is non-constant, and
	assign size with indirect reference.  Reorder elements added to
	iterator vector and add element containing the iteration count.  Use
	omp_iterator_elems_length to compute element array size required.
	* gimplify.h (omp_iterator_elems_length): New prototype.
	* omp-low.cc (lower_omp_map_iterator_expr): Reorder elements read
	from iterator vector.  If elements field is a pointer type, assign
	using pointer arithmetic followed by indirect reference, and return
	the field directly.
	(lower_omp_map_iterator_size): Reorder elements read from iterator
	vector.  If elements field is a pointer type, assign using pointer
	arithmetic followed by indirect reference.
	(allocate_omp_iterator_elems): New.
	(free_omp_iterator_elems): New.
	(lower_omp_target): Call allocate_omp_iterator_elems before inserting
	loops sequence, and call free_omp_iterator_elems afterwards.
	* tree-pretty-print.cc (dump_omp_iterators): Print extra elements in
	iterator vector.

gcc/testsuite/

	* c-c++-common/gomp/target-map-iterators-3.c: Update expected Gimple
	output.
	* c-c++-common/gomp/target-map-iterators-5.c: New.
	* c-c++-common/gomp/target-update-iterators-3.c: Update expected
	Gimple output.
	* gfortran.dg/gomp/target-map-iterators-3.f90: Likewise.
	* gfortran.dg/gomp/target-map-iterators-5.f90: New.
	* gfortran.dg/gomp/target-update-iterators-3.f90: Update expected
	Gimple output.

libgomp/

	* testsuite/libgomp.c-c++-common/target-map-iterators-4.c: New.
	* testsuite/libgomp.c-c++-common/target-map-iterators-5.c: New.
	* testsuite/libgomp.c-c++-common/target-update-iterators-4.c: New.
	* testsuite/libgomp.fortran/target-map-iterators-4.f90: New.
	* testsuite/libgomp.fortran/target-map-iterators-5.f90: New.
	* testsuite/libgomp.fortran/target-update-iterators-4.f90: New.

10 files changed, 166 insertions, 46 deletions

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 9164875..5adec95 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10068,6 +10068,13 @@ struct iterator_loop_info_t
 
 typedef hash_map<tree, iterator_loop_info_t> iterator_loop_info_map_t;
 
+tree
+omp_iterator_elems_length (tree count)
+{
+  tree count_2 = size_binop (MULT_EXPR, count, size_int (2));
+  return size_binop (PLUS_EXPR, count_2, size_int (1));
+}
+
 /* Builds a loop to expand any OpenMP iterators in the clauses in LIST_P,
    reusing any previously built loops if they use the same set of iterators.
    Generated Gimple statements are placed into LOOPS_SEQ_P.  The clause
@@ -10123,27 +10130,21 @@ build_omp_iterators_loops (tree *list_p, gimple_seq *loops_seq_p)
 	}
 
       /* 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, "omp_iter_data");
-	  TREE_ADDRESSABLE (elems) = 1;
-	  gimple_add_tmp_var (elems);
-	}
-      else
-	{
-	  /* Handle dynamic sizes.  */
-	  sorry ("dynamic iterator sizes not implemented yet");
-	}
+      tree arr_length = omp_iterator_elems_length (loop.count);
+      tree elems_type = TREE_CONSTANT (arr_length)
+		? build_array_type (ptr_type_node,
+				    build_index_type (arr_length))
+		: build_pointer_type (ptr_type_node);
+      tree elems = create_tmp_var_raw (elems_type, "omp_iter_data");
+      TREE_ADDRESSABLE (elems) = 1;
+      gimple_add_tmp_var (elems);
 
       /* BEFORE LOOP:  */
       /* elems[0] = count;  */
-      tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
-			 NULL_TREE, NULL_TREE);
+      tree lhs = TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE
+	? build4 (ARRAY_REF, ptr_type_node, elems, size_int (0), NULL_TREE,
+		  NULL_TREE)
+	: build1 (INDIRECT_REF, ptr_type_node, elems);
       tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
 			     void_type_node, lhs, loop.count);
       gimplify_and_add (tem, loops_seq_p);
@@ -10151,10 +10152,11 @@ build_omp_iterators_loops (tree *list_p, gimple_seq *loops_seq_p)
       /* Make a copy of the iterator with extra info at the end.  */
       int elem_count = TREE_VEC_LENGTH (OMP_CLAUSE_ITERATORS (c));
       tree new_iterator = copy_omp_iterator (OMP_CLAUSE_ITERATORS (c),
-					     elem_count + 3);
+					     elem_count + 4);
       TREE_VEC_ELT (new_iterator, elem_count) = loop.body_label;
-      TREE_VEC_ELT (new_iterator, elem_count + 1) = elems;
-      TREE_VEC_ELT (new_iterator, elem_count + 2) = loop.index;
+      TREE_VEC_ELT (new_iterator, elem_count + 1) = loop.index;
+      TREE_VEC_ELT (new_iterator, elem_count + 2) = elems;
+      TREE_VEC_ELT (new_iterator, elem_count + 3) = loop.count;
       TREE_CHAIN (new_iterator) = TREE_CHAIN (OMP_CLAUSE_ITERATORS (c));
       OMP_CLAUSE_ITERATORS (c) = new_iterator;
 
diff --git a/gcc/gimplify.h b/gcc/gimplify.h
index 80c335e..ca970cf 100644
--- a/gcc/gimplify.h
+++ b/gcc/gimplify.h
@@ -79,6 +79,7 @@ extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *,
 extern tree omp_get_construct_context (void);
 int omp_has_novariants (void);
 
+extern tree omp_iterator_elems_length (tree count);
 extern gimple_seq *enter_omp_iterator_loop_context (tree, gomp_target *,
 						    gimple_seq * = NULL);
 extern void exit_omp_iterator_loop_context (tree);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 48e78ef..8fcae3d 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -13601,20 +13601,30 @@ lower_omp_map_iterator_expr (tree expr, tree c, gomp_target *stmt)
     return expr;
 
   tree iterator = OMP_CLAUSE_ITERATORS (c);
-  tree elems = TREE_VEC_ELT (iterator, 7);
-  tree index = TREE_VEC_ELT (iterator, 8);
+  tree index = TREE_VEC_ELT (iterator, 7);
+  tree elems = TREE_VEC_ELT (iterator, 8);
   gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt);
 
    /* IN LOOP BODY:  */
    /* elems[idx] = <expr>;  */
-  tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, index,
-		     NULL_TREE, NULL_TREE);
+  tree lhs;
+  if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+    lhs = build4 (ARRAY_REF, ptr_type_node, elems, index, NULL_TREE, NULL_TREE);
+  else
+    {
+      tree tmp = size_binop (MULT_EXPR, index, TYPE_SIZE_UNIT (ptr_type_node));
+      tmp = size_binop (POINTER_PLUS_EXPR, elems, tmp);
+      lhs = build1 (INDIRECT_REF, ptr_type_node, tmp);
+    }
   tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
 			      void_type_node, lhs, expr);
   gimplify_and_add (mod_expr, loop_body_p);
   exit_omp_iterator_loop_context (c);
 
-  return build_fold_addr_expr_with_type (elems, ptr_type_node);
+  if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+    return build_fold_addr_expr_with_type (elems, ptr_type_node);
+  else
+    return elems;
 }
 
 /* Set SIZE as the size expression that should result from the clause C.
@@ -13628,15 +13638,25 @@ lower_omp_map_iterator_size (tree size, tree c, gomp_target *stmt)
     return size;
 
   tree iterator = OMP_CLAUSE_ITERATORS (c);
-  tree elems = TREE_VEC_ELT (iterator, 7);
-  tree index = TREE_VEC_ELT (iterator, 8);
+  tree index = TREE_VEC_ELT (iterator, 7);
+  tree elems = TREE_VEC_ELT (iterator, 8);
   gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt);
 
   /* IN LOOP BODY:  */
   /* elems[idx+1] = <size>;  */
-  tree lhs = build4 (ARRAY_REF, ptr_type_node, elems,
-		     size_binop (PLUS_EXPR, index, size_int (1)),
-		     NULL_TREE, NULL_TREE);
+  tree lhs;
+  if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+    lhs = build4 (ARRAY_REF, ptr_type_node, elems,
+		  size_binop (PLUS_EXPR, index, size_int (1)),
+		  NULL_TREE, NULL_TREE);
+  else
+    {
+      tree index_1 = size_binop (PLUS_EXPR, index, size_int (1));
+      tree tmp = size_binop (MULT_EXPR, index_1,
+			     TYPE_SIZE_UNIT (ptr_type_node));
+      tmp = size_binop (POINTER_PLUS_EXPR, elems, tmp);
+      lhs = build1 (INDIRECT_REF, ptr_type_node, tmp);
+    }
   tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
 			      void_type_node, lhs, size);
   gimplify_and_add (mod_expr, loop_body_p);
@@ -13645,6 +13665,62 @@ lower_omp_map_iterator_size (tree size, tree c, gomp_target *stmt)
   return size_int (SIZE_MAX);
 }
 
+static void
+allocate_omp_iterator_elems (tree clauses, gimple_seq loops_seq)
+{
+  for (tree c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (!OMP_CLAUSE_HAS_ITERATORS (c))
+	continue;
+      tree iters = OMP_CLAUSE_ITERATORS (c);
+      tree elems = TREE_VEC_ELT (iters, 8);
+      if (!POINTER_TYPE_P (TREE_TYPE (elems)))
+	continue;
+      tree arr_length = omp_iterator_elems_length (TREE_VEC_ELT (iters, 9));
+      tree call = builtin_decl_explicit (BUILT_IN_MALLOC);
+      tree size = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MULT_EXPR,
+				   size_type_node, arr_length,
+				   TYPE_SIZE_UNIT (ptr_type_node));
+      tree tmp = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), call, 1,
+				      size);
+
+      /* Find the first statement '<index> = -1' in the pre-loop statements.  */
+      tree index = TREE_VEC_ELT (iters, 7);
+      gimple_stmt_iterator gsi;
+      for (gsi = gsi_start (loops_seq); !gsi_end_p (gsi); gsi_next (&gsi))
+	{
+	  gimple *stmt = gsi_stmt (gsi);
+	  if (gimple_code (stmt) == GIMPLE_ASSIGN
+	      && gimple_assign_lhs (stmt) == index
+	      && gimple_assign_rhs1 (stmt) == size_int (-1))
+	    break;
+	}
+      gcc_assert (!gsi_end_p (gsi));
+
+      gimple_seq alloc_seq = NULL;
+      gimplify_assign (elems, tmp, &alloc_seq);
+      gsi_insert_seq_before (&gsi, alloc_seq, GSI_SAME_STMT);
+    }
+}
+
+static void
+free_omp_iterator_elems (tree clauses, gimple_seq *seq)
+{
+  for (tree c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (!OMP_CLAUSE_HAS_ITERATORS (c))
+	continue;
+      tree elems = TREE_VEC_ELT (OMP_CLAUSE_ITERATORS (c), 8);
+      if (!POINTER_TYPE_P (TREE_TYPE (elems)))
+	continue;
+      tree call = builtin_decl_explicit (BUILT_IN_FREE);
+      call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), call, 1, elems);
+      gimplify_and_add (call, seq);
+      tree clobber = build_clobber (TREE_TYPE (elems));
+      gimple_seq_add_stmt (seq, gimple_build_assign (elems, clobber));
+    }
+}
+
 /* Lower the GIMPLE_OMP_TARGET in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
@@ -15810,9 +15886,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gimple_omp_set_body (stmt, new_body);
     }
 
+  allocate_omp_iterator_elems (clauses,
+			       gimple_omp_target_iterator_loops (stmt));
   gsi_insert_seq_before (gsi_p, gimple_omp_target_iterator_loops (stmt),
 			 GSI_SAME_STMT);
   gimple_omp_target_set_iterator_loops (stmt, NULL);
+  free_omp_iterator_elems (clauses, &olist);
+
   bind = gimple_build_bind (NULL, NULL,
 			    tgt_bind ? gimple_bind_block (tgt_bind)
 				     : NULL_TREE);
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
index 62df42f..87b32e4 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
@@ -17,7 +17,7 @@ void f (int ***x, float ***y, double **z)
 
 /* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 3 "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, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):from:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):to:\\*D\\\.\[0-9\]+" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):from:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):to:\\*D\\\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 4 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
new file mode 100644
index 0000000..e79e6a5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#define DIM2 17
+
+void f (int **x, int lbound, int ubound, int stride)
+{
+  #pragma omp target map(to:x) map(iterator(i=lbound:ubound:stride), to: x[i][:DIM2])
+    ;
+}
+
+/* { dg-final { scan-tree-dump-times "_\[0-9\]+ = ubound - lbound;" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "D\\\.\[0-9\]+ = __builtin_malloc \\(D\\\.\[0-9\]+\\);" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_free \\(omp_iter_data\\\.\[0-9\]+\\);" 2 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
index ef55216..6618962 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
@@ -13,5 +13,5 @@ void f (int ***x, float ***y, double **z)
 
 /* { 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, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+, index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+, index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\.\[0-9\]+>, index=D\.\[0-9\]+, elems=omp_iter_data\.\[0-9\]+, elems_count=200\\):\\*D\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1, loop_label=<D\.\[0-9\]+>, index=D\.\[0-9\]+, elems=omp_iter_data\.\[0-9\]+, elems_count=10\\):\\*D\.\[0-9\]+" 1 "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
index 7dad2a6..1099955 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
@@ -18,7 +18,7 @@ end program
 
 ! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
 ! { dg-final { scan-tree-dump-times "if \\(i <= 27\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):to:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):from:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):attach:x\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):attach:y\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):to:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=27\\):from:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):attach:x\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=27\\):attach:y\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
new file mode 100644
index 0000000..f620d1c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+module m
+  integer, parameter :: DIM1 = 31
+  integer, parameter :: DIM2 = 17
+  type :: array_ptr
+    integer, pointer :: ptr(:)
+ end type
+contains
+  subroutine f (x, stride)
+    type (array_ptr) :: x(:)
+    integer :: stride
+
+    !$omp target map(to: x) map(iterator(i=lbound(x, 1):ubound(x, 1):stride), to: x(i)%ptr(:))
+    !$omp end target
+  end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]+ = __builtin_malloc \\(D\\\.\[0-9\]+\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_free \\(omp_iter_data\\\.\[0-9\]+\\);" 3 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
index d9c92cf..a8dffcf 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
@@ -18,6 +18,6 @@ program test
 end program
 
 ! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "if \\(j <= 39\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\) j=1:39:1, integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "from\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump "if \\(j <= 39\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" "gimple" } }
+! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\) j=1:39:1, integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=663\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 2 "gimple" } }
+! { dg-final { scan-tree-dump "from\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" "gimple" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 79e212c..cb96859 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -452,10 +452,12 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags)
     {
       pp_string (pp, ", loop_label=");
       dump_generic_node (pp, TREE_VEC_ELT (iter, 6), spc, flags, false);
-      pp_string (pp, ", elems=");
-      dump_generic_node (pp, TREE_VEC_ELT (iter, 7), spc, flags, false);
       pp_string (pp, ", index=");
+      dump_generic_node (pp, TREE_VEC_ELT (iter, 7), spc, flags, false);
+      pp_string (pp, ", elems=");
       dump_generic_node (pp, TREE_VEC_ELT (iter, 8), spc, flags, false);
+      pp_string (pp, ", elems_count=");
+      dump_generic_node (pp, TREE_VEC_ELT (iter, 9), spc, flags, false);
     }
   pp_right_paren (pp);
 }