openmp: Add support for iterators in 'target update' clauses (C/C++)

This adds support for iterators in 'to' and 'from' clauses in the
'target update' OpenMP directive.

gcc/c/

	* c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier.
	* c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from
	clauses.

gcc/cp/

	* parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier.
	* semantics.cc (finish_omp_clauses): Finish iterators for to/from
	clauses.

gcc/

	* gimplify.cc (remove_unused_omp_iterator_vars): Display unused
	variable warning for 'to' and 'from' clauses.
	(gimplify_scan_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): Add argument for iterator loops sequence
	in call to gimplify_scan_omp_clauses.
	(gimplify_omp_target_update): Call remove_unused_omp_iterator_vars and
	build_omp_iterators_loops.  Add loop sequence as argument when calling
	gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building
	the Gimple statement.
	* tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators
	for to/from clauses with iterators.
	* tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM
	and OMP_CLAUSE_TO.
	* tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and
	OMP_CLAUSE_FROM.
	(OMP_CLAUSE_ITERATORS): Likewise.

gcc/testsuite/

	* c-c++-common/gomp/target-update-iterators-1.c: New.
	* c-c++-common/gomp/target-update-iterators-2.c: New.
	* c-c++-common/gomp/target-update-iterators-3.c: New.

libgomp/

	* target.c (gomp_update): Call gomp_merge_iterator_maps.  Free
	allocated variables.
	* testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New.
	* testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New.
	* testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New.

3 files changed, 190 insertions, 0 deletions

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 0000000..5a4cad5
--- /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 0000000..93438d0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
@@ -0,0 +1,58 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 0000000..a70b21c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
@@ -0,0 +1,67 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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;
+}