openmp: Add support for using custom mappers with iterators (C, C++)

gcc/c-family/

	* c-omp.cc (omp_instantiate_mapper): Apply iterator to new clauses
	generated from mapper.

gcc/c/

	* c-parser.cc (c_parser_omp_clause_map): Apply iterator to push and
	pop mapper clauses.

gcc/cp/

	* parser.cc (cp_parser_omp_clause_map): Apply iterator to push and
	pop mapper clauses.
	* semantics.cc (cxx_omp_map_array_section): Allow array types for
	base type of array sections.

libgomp/

	* testsuite/libgomp.c-c++-common/mapper-iterators-1.c: New test.
	* testsuite/libgomp.c-c++-common/mapper-iterators-2.c: New test.
	* testsuite/libgomp.c-c++-common/mapper-iterators-3.c: New test.

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

7 files changed, 274 insertions, 1 deletions

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 7d33318..a3cd1ae 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4761,6 +4761,7 @@ omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
   tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
   tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
   tree mapper_name = NULL_TREE;
+  tree iterator = *outlist ? OMP_CLAUSE_ITERATORS (*outlist) : NULL_TREE;
 
   remap_mapper_decl_info map_info;
   map_info.dummy_var = dummy_var;
@@ -4889,6 +4890,7 @@ omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
 	}
       else
 	{
+	  OMP_CLAUSE_ITERATORS (unshared) = iterator;
 	  *outlist = unshared;
 	  outlist = &OMP_CLAUSE_CHAIN (unshared);
 	}
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 8a5b493..31ab7d0 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20738,6 +20738,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
       tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
       OMP_CLAUSE_DECL (name) = mapper_name;
+      if (iterators)
+	OMP_CLAUSE_ITERATORS (name) = iterators;
       OMP_CLAUSE_CHAIN (name) = nl;
       nl = name;
 
@@ -20746,6 +20748,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
       name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
       OMP_CLAUSE_DECL (name) = null_pointer_node;
+      if (iterators)
+	OMP_CLAUSE_ITERATORS (name) = iterators;
       OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
       OMP_CLAUSE_CHAIN (last_new) = name;
     }
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 83e0220..f38a327 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -43344,6 +43344,8 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
       name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
       OMP_CLAUSE_DECL (name) = null_pointer_node;
+      if (iterators)
+	OMP_CLAUSE_ITERATORS (name) = iterators;
       OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
       OMP_CLAUSE_CHAIN (last_new) = name;
     }
@@ -43659,6 +43661,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind)
       tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
       OMP_CLAUSE_DECL (name) = mapper_name;
+      if (iterators)
+	OMP_CLAUSE_ITERATORS (name) = iterators;
       OMP_CLAUSE_CHAIN (name) = nlist;
       nlist = name;
 
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index b80bf33..770872a 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -6998,7 +6998,8 @@ cxx_omp_map_array_section (location_t loc, tree t)
       if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
 	t = convert_from_reference (t);
 
-      if (TYPE_PTR_P (TREE_TYPE (t)))
+      if (TYPE_PTR_P (TREE_TYPE (t))
+	  || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
 	t = build_array_ref (loc, t, low);
       else
 	t = error_mark_node;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-1.c
new file mode 100644
index 0000000..1938237
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-1.c
@@ -0,0 +1,83 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdio.h>
+#include <assert.h>
+
+#define DIM1 4
+#define DIM2 16
+
+struct S {
+  int *arr1;
+  float *arr2;
+  size_t len;
+};
+
+size_t
+mkarray (struct S arr[])
+{
+  size_t sum = 0;
+
+  for (int i = 0; i < DIM1; i++)
+    {
+      memset (&arr[i], 0, sizeof (struct S));
+      arr[i].len = DIM2;
+      arr[i].arr1 = (int *) calloc (arr[i].len, sizeof (int));
+      for (int j = 0; j < DIM2; j++)
+	{
+	  size_t value = (i + 1) * (j + 1);
+	  sum += value;
+	  arr[i].arr1[j] = value;
+	}
+    }
+
+  return sum;
+}
+
+int main ()
+{
+  struct S arr[DIM1];
+  size_t sum = 0xdeadbeef;
+  size_t expected = mkarray (arr);
+
+  #pragma omp declare mapper (struct S x) \
+	map(to: x.arr1[0:DIM2]) \
+	map(to: x.arr2[0:DIM2]) \
+	map(to: x.len)
+
+  #pragma omp target map(iterator(int i=0:DIM1), to: arr[i]) map(from: sum)
+    {
+      sum = 0;
+#ifdef DEBUG
+      __builtin_printf ("&sum: %p\n", &sum);
+#endif
+      for (int i = 0; i < DIM1; i++)
+	{
+#ifdef DEBUG
+	  __builtin_printf ("&arr[%d] = %p\n", i, &arr[i]);
+	  __builtin_printf ("arr[%d].len = %d\n", i, arr[i].len);
+	  __builtin_printf ("arr[%d].arr1 = %p\n", i, arr[i].arr1);
+	  __builtin_printf ("arr[%d].arr2 = %p\n", i, arr[i].arr2);
+#endif
+	  for (int j = 0; j < DIM2; j++)
+	    {
+#ifdef DEBUG
+	      __builtin_printf ("(i=%d,j=%d): %p\n", i, j, &arr[i].arr1[j]);
+	      __builtin_printf ("(i=%d,j=%d): %d\n", i, j, arr[i].arr1[j]);
+#endif
+	      sum += arr[i].arr1[j];
+#ifdef DEBUG
+	      __builtin_printf ("sum: %ld\n", sum);
+#endif
+	    }
+	}
+    }
+
+#ifdef DEBUG
+  __builtin_printf ("&sum: %p\n", &sum);
+  __builtin_printf ("sum:%zd (expected: %zd)\n", sum, expected);
+#endif
+
+  return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-2.c
new file mode 100644
index 0000000..76f00fb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-2.c
@@ -0,0 +1,81 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdio.h>
+#include <assert.h>
+
+#define DIM1 4
+#define DIM2 16
+
+#ifdef DEBUG
+#undef DEBUG
+#define DEBUG(...) __builtin_printf (__VA_ARGS__)
+#else
+#define DEBUG(...)
+#endif
+
+struct S {
+  int *arr1;
+  float *arr2;
+  size_t len;
+};
+
+size_t
+mkarray (struct S arr[])
+{
+  size_t sum = 0;
+
+  for (int i = 0; i < DIM1; i++)
+    {
+      memset (&arr[i], 0, sizeof (struct S));
+      arr[i].len = DIM2;
+      arr[i].arr1 = (int *) calloc (arr[i].len, sizeof (int));
+      for (int j = 0; j < DIM2; j++)
+	{
+	  size_t value = (i + 1) * (j + 1);
+	  sum += value;
+	  arr[i].arr1[j] = value;
+	}
+    }
+
+  return sum;
+}
+
+int main ()
+{
+  struct S arr[DIM1];
+  size_t sum = 0xdeadbeef;
+  size_t expected = mkarray (arr);
+
+  #pragma omp declare mapper (struct S x) \
+	map(to: x.arr1[0:DIM2]) \
+	map(to: x.arr2[0:DIM2]) \
+	map(to: x.len)
+
+  /* This should be equivalent to map(iterator(int i=0:DIM1), to: arr[i])  */
+  #pragma omp target map(iterator(int i=0:DIM1:2, j=0:2), to: arr[i+j]) map(from: sum)
+    {
+      sum = 0;
+      DEBUG ("&sum: %p\n", &sum);
+      for (int i = 0; i < DIM1; i++)
+	{
+	  DEBUG ("&arr[%d] = %p\n", i, &arr[i]);
+	  DEBUG ("arr[%d].len = %d\n", i, arr[i].len);
+	  DEBUG ("arr[%d].arr1 = %p\n", i, arr[i].arr1);
+	  DEBUG ("arr[%d].arr2 = %p\n", i, arr[i].arr2);
+	  for (int j = 0; j < DIM2; j++)
+	    {
+	      DEBUG ("(i=%d,j=%d): %p\n", i, j, &arr[i].arr1[j]);
+	      DEBUG ("(i=%d,j=%d): %d\n", i, j, arr[i].arr1[j]);
+	      sum += arr[i].arr1[j];
+	      DEBUG ("sum: %ld\n", sum);
+	    }
+	}
+    }
+
+  DEBUG ("&sum: %p\n", &sum);
+  DEBUG ("sum:%zd (expected: %zd)\n", sum, expected);
+
+  return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-3.c
new file mode 100644
index 0000000..9d67c38
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-3.c
@@ -0,0 +1,98 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+#define DIM 4
+
+#ifdef DEBUG
+#undef DEBUG
+#define DEBUG(...) __builtin_printf (__VA_ARGS__)
+#else
+#define DEBUG(...)
+#endif
+
+typedef struct {
+  int *arr;
+  int size;
+} B;
+
+#pragma omp declare mapper (mapB : B myb) map(to: myb.size, myb.arr) \
+					  map(tofrom: myb.arr[0:myb.size+1])
+
+struct A {
+  int *arr1;
+  B *arr2;
+  int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct A var[DIM];
+
+  for (int i=0; i < DIM; i++)
+    {
+      memset (&var[i], 0, sizeof var[i]);
+      var[i].arr1 = (int *) calloc (N, sizeof (int));
+      var[i].arr2 = (B *) malloc (sizeof (B));
+      var[i].arr2->arr = (int *) calloc (N+1, sizeof (float));
+      var[i].arr2->size = N+1;
+      DEBUG ("host &var[%d]:%p\n", i, &var[i]);
+      DEBUG ("host var[%d].arr1:%p\n", i, var[i].arr1);
+      DEBUG ("host var[%d].arr2:%p\n", i, var[i].arr2);
+      DEBUG ("host var[%d].arr2->arr:%p\n", i, var[i].arr2->arr);
+      DEBUG ("host var[%d].arr2->size:%d\n", i, var[i].arr2->size);
+    }
+
+  {
+    #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+			  map(tofrom: x.arr1[0:N]) \
+			  map(mapper(mapB), tofrom: x.arr2[0:1])
+    #pragma omp target map(iterator(int i=0:DIM), tofrom: var[i])
+    {
+      for (int i = 0; i < DIM; i++)
+	{
+	  DEBUG ("&var[%d]:%p\n", i, &var[i]);
+	  DEBUG ("var[%d].arr1:%p\n", i, var[i].arr1);
+	  DEBUG ("var[%d].arr2:%p\n", i, var[i].arr2);
+	  if (var[i].arr2)
+	    {
+	      DEBUG ("var[%d].arr2->arr:%p\n", i, var[i].arr2->arr);
+	      DEBUG ("var[%d].arr2->size:%d\n", i, var[i].arr2->size);
+	    }
+	  for (int j = 0; j < N; j++)
+	    {
+	      DEBUG ("&var[%d].arr1[%d]:%p\n", i, j, &var[i].arr1[j]);
+	      var[i].arr1[j]++;
+	      if (var[i].arr2)
+		{
+		  DEBUG ("&var[%d].arr2->arr[%d]:%p\n", i, j, &var[i].arr2->arr[j]);
+		  var[i].arr2->arr[j]++;
+		}
+	      else
+		DEBUG ("SKIP arr2\n");
+	    }
+	}
+    }
+  }
+
+  for (int i = 0; i < DIM; i++)
+    for (int j = 0; j < N; j++)
+      {
+	assert (var[i].arr1[j] == 1);
+	assert (var[i].arr2->arr[j] == 1);
+	assert (var[i].arr3[j] == 0);
+      }
+
+  for (int i = 0; i < DIM; i++)
+    {
+      free (var[i].arr1);
+      free (var[i].arr2->arr);
+      free (var[i].arr2);
+    }
+
+  return 0;
+}