[gcc/devel/omp/gcc-16] openmp: Add support for using custom mappers with iterators (C, C++)

https://gcc.gnu.org/g:4afbd256410d6050d44a213d1108ed724a4d1405

commit 4afbd256410d6050d44a213d1108ed724a4d1405
Author: Kwok Cheung Yeung <[email protected]>
Date:   Mon Jan 13 13:08:07 2025 +0000

    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.
    
    libgomp/
    
            * testsuite/libgomp.c-c++-common/mapper-iterators-1.c: New test.
            * testsuite/libgomp.c-c++-common/mapper-iterators-2.c: New test.
    
    Co-authored-by: Andrew Stubbs <[email protected]>

Diff:
---
 gcc/c-family/c-omp.cc                              |  2 +
 gcc/c/c-parser.cc                                  |  4 ++
 gcc/cp/parser.cc                                   |  4 ++
 .../libgomp.c-c++-common/mapper-iterators-1.c      | 83 ++++++++++++++++++++++
 .../libgomp.c-c++-common/mapper-iterators-2.c      | 81 +++++++++++++++++++++
 5 files changed, 174 insertions(+)

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index e263706a3f1e..76845f01cdc9 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4603,6 +4603,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;
@@ -4732,6 +4733,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 e8077b216dfa..08a1b1e005c3 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20970,6 +20970,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list, bool declare_mapper_p)
       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;
 
@@ -20978,6 +20980,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list, bool declare_mapper_p)
       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 0418d44172d4..9f305723acee 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45741,6 +45741,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, bool declare_mapper_p)
       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;
 
@@ -45749,6 +45751,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, bool declare_mapper_p)
       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/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 000000000000..193823744bd6
--- /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 000000000000..76f00fb518a6
--- /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;
+}