openmp: Add support for iterators in map clauses (C/C++)

This adds preliminary support for iterators in map clauses within OpenMP
'target' constructs (which includes constructs such as 'target enter data').

Iterators with non-constant loop bounds are not currently supported.

gcc/c/

	* c-parser.cc (c_parser_omp_variable_list): Use location of the
	map expression as the clause location.
	(c_parser_omp_clause_map): Parse 'iterator' modifier.
	* c-typeck.cc (c_finish_omp_clauses): Finish iterators.  Apply
	iterators to generated clauses.

gcc/cp/

	* parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier.
	* semantics.cc (finish_omp_clauses): Finish iterators.  Apply
	iterators to generated clauses.

gcc/

	* gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded
	iterator loops.
	* gimple.cc (gimple_build_omp_target): Add argument for iterator
	loops sequence.  Initialize iterator loops field.
	* gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET.
	* gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET.  Add extra
	field for iterator loops.
	(gimple_build_omp_target): Add argument for iterator loops sequence.
	(gimple_omp_target_iterator_loops): New.
	(gimple_omp_target_iterator_loops_ptr): New.
	(gimple_omp_target_set_iterator_loops): New.
	* gimplify.cc (find_var_decl): New.
	(copy_omp_iterator): New.
	(remap_omp_iterator_var_1): New.
	(remap_omp_iterator_var): New.
	(remove_unused_omp_iterator_vars): New.
	(struct iterator_loop_info_t): New type.
	(iterator_loop_info_map_t): New type.
	(build_omp_iterators_loops): New.
	(enter_omp_iterator_loop_context_1): New.
	(enter_omp_iterator_loop_context): New.
	(enter_omp_iterator_loop_context): New.
	(exit_omp_iterator_loop_context): New.
	(gimplify_adjust_omp_clauses): Add argument for iterator loop
	sequence.  Gimplify the clause decl and size into the iterator
	loop if iterators are used.
	(gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and
	build_omp_iterators_loops for OpenMP target expressions.  Add
	loop sequence as argument when calling gimplify_adjust_omp_clauses
	and building the Gimple statement.
	* gimplify.h (enter_omp_iterator_loop_context): New prototype.
	(exit_omp_iterator_loop_context): New prototype.
	* gsstruct.def (GSS_OMP_TARGET): New.
	* omp-low.cc (lower_omp_map_iterator_expr): New.
	(lower_omp_map_iterator_size): New.
	(finish_omp_map_iterators): New.
	(lower_omp_target): Add sorry if iterators used with deep mapping.
	Call lower_omp_map_iterator_expr before assigning to sender ref.
	Call lower_omp_map_iterator_size before setting the size.  Insert
	iterator loop sequence before the statements for the target clause.
	* tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator
	loop sequence of OpenMP target statements.
	(convert_local_reference_stmt): Likewise.
	(convert_tramp_reference_stmt): Likewise.
	* tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator
	information if present.
	(dump_omp_clause): Call dump_omp_iterators for iterators in map
	clauses.
	* tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP.
	(walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP.
	* tree.h (OMP_CLAUSE_HAS_ITERATORS): New.
	(OMP_CLAUSE_ITERATORS): New.

gcc/testsuite/

	* c-c++-common/gomp/map-6.c (foo): Amend expected error message.
	* c-c++-common/gomp/target-map-iterators-1.c: New.
	* c-c++-common/gomp/target-map-iterators-2.c: New.
	* c-c++-common/gomp/target-map-iterators-3.c: New.
	* c-c++-common/gomp/target-map-iterators-4.c: New.

libgomp/

	* target.c (kind_to_name): New.
	(gomp_merge_iterator_maps): New.
	(gomp_map_vars_internal): Call gomp_merge_iterator_maps.  Copy
	address of only the first iteration to target vars.  Free allocated
	variables.
	* testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New.
	* testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New.
	* testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New.

Co-authored-by: Andrew Stubbs <ams@baylibre.com>

3 files changed, 147 insertions, 0 deletions

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 0000000..b3d87f2
--- /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 0000000..8569b55
--- /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 0000000..be30fa65d
--- /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;
+}