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>

4 files changed, 325 insertions, 7 deletions

diff --git a/libgomp/target.c b/libgomp/target.c
index cda092b..43976f0 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1003,6 +1003,155 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
     }
 }
 
+static const char *
+kind_to_name (unsigned short kind, bool short_mapkind)
+{
+  if (short_mapkind && GOMP_MAP_IMPLICIT_P (kind))
+    kind &= ~GOMP_MAP_IMPLICIT;
+
+  switch (kind & (short_mapkind ? 0xff : 0x7))
+    {
+    case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+    case GOMP_MAP_TO: return "GOMP_MAP_TO";
+    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_TO_PSET: return "GOMP_MAP_TO_PSET";
+    case GOMP_MAP_FORCE_PRESENT: return "GOMP_MAP_FORCE_PRESENT";
+    case GOMP_MAP_DELETE: return "GOMP_MAP_DELETE";
+    case GOMP_MAP_FORCE_DEVICEPTR: return "GOMP_MAP_FORCE_DEVICEPTR";
+    case GOMP_MAP_DEVICE_RESIDENT: return "GOMP_MAP_DEVICE_RESIDENT";
+    case GOMP_MAP_LINK: return "GOMP_MAP_LINK";
+    case GOMP_MAP_IF_PRESENT: return "GOMP_MAP_IF_PRESENT";
+    case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+    case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+    case GOMP_MAP_USE_DEVICE_PTR: return "GOMP_MAP_USE_DEVICE_PTR";
+    case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: return "GOMP_MAP_ZERO_LEN_ARRAY_SECTION";
+    case GOMP_MAP_FORCE_ALLOC: return "GOMP_MAP_FORCE_ALLOC";
+    case GOMP_MAP_FORCE_TO: return "GOMP_MAP_FORCE_TO";
+    case GOMP_MAP_FORCE_FROM: return "GOMP_MAP_FORCE_FROM";
+    case GOMP_MAP_FORCE_TOFROM: return "GOMP_MAP_FORCE_TOFROM";
+    case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
+      return "GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT";
+    case GOMP_MAP_ALWAYS_TO: return "GOMP_MAP_ALWAYS_TO";
+    case GOMP_MAP_ALWAYS_FROM: return "GOMP_MAP_ALWAYS_FROM";
+    case GOMP_MAP_ALWAYS_TOFROM: return "GOMP_MAP_ALWAYS_TOFROM";
+    case GOMP_MAP_ALWAYS_PRESENT_TO: return "GOMP_MAP_ALWAYS_PRESENT_TO";
+    case GOMP_MAP_ALWAYS_PRESENT_FROM: return "GOMP_MAP_ALWAYS_PRESENT_FROM";
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM: return "GOMP_MAP_ALWAYS_PRESENT_TOFROM";
+    case GOMP_MAP_STRUCT: return "GOMP_MAP_STRUCT";
+    case GOMP_MAP_STRUCT_UNORD: return "GOMP_MAP_STRUCT_UNORD";
+    case GOMP_MAP_ALWAYS_POINTER: return "GOMP_MAP_ALWAYS_POINTER";
+    case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+      return "GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION";
+    case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
+      return "GOMP_MAP_DELETE_ZERO_LENGTH_ARRAY_SECTION";
+    case GOMP_MAP_RELEASE: return "GOMP_MAP_RELEASE";
+    case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+    case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+    case GOMP_MAP_FORCE_DETACH: return "GOMP_MAP_FORCE_DETACH";
+    case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+      return "GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION";
+    default: return "unknown";
+    }
+}
+
+/*  When GCC encounters a clause with an iterator, e.g.:
+
+   #pragma omp target map (iterator(i=0:4), to: x[i])
+
+   it generates an array containing the number of iterations and the
+   address and size of each iteration. e.g.:
+
+   void *omp_iter_data[] = {
+     (void *) 4, // Number of iterations
+     &x[0], (void *) sizeof(x[0]),
+     &x[1], (void *) sizeof(x[1]),
+     &x[2], (void *) sizeof(x[2]),
+     &x[3], (void *) sizeof(x[3])
+   };
+
+   When the construct is lowered, &omp_iter_data is used as the host address
+   for the map (instead of &x[i]), and the size is set to SIZE_MAX to mark
+   the map as an iterator map.
+
+   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 (and 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], true),
+			  (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,
@@ -1019,6 +1168,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)  /* OpenMP */
+    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;
@@ -1896,14 +2050,22 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 
   if (pragma_kind & GOMP_MAP_VARS_TARGET)
     {
+      /* The target variables table is constructed with maps using iterators
+	 unexpanded. Now that the iterator maps are expanded, we will need to
+	 skip all expanded maps after the initial entry, otherwise subsequent
+	 maps will be out-of-sync with their corresponding entry in the
+	 target variables table.  */
+      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)
@@ -1935,6 +2097,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 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;
+}