From: Kwok Cheung Yeung <[email protected]>
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.
* testsuite/libgomp.c-c++-common/iterator-with-structs.c: New test.
Co-authored-by: Andrew Stubbs <[email protected]>
Co-authored-by: Sandra Loosemore <[email protected]>
---
gcc/c-family/c-omp.cc | 2 +
gcc/c/c-parser.cc | 4 +
gcc/cp/parser.cc | 4 +
.../iterator-with-structs.c | 24 ++++++
.../libgomp.c-c++-common/mapper-iterators-1.c | 83 +++++++++++++++++++
.../libgomp.c-c++-common/mapper-iterators-2.c | 81 ++++++++++++++++++
6 files changed, 198 insertions(+)
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/iterator-with-structs.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/mapper-iterators-2.c
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 181dfb3c7b7..6e62793363d 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4601,6 +4601,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;
@@ -4730,6 +4731,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 af9d48b89bf..541d4c91bae 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20994,6 +20994,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;
@@ -21002,6 +21004,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 2ab4d7d968f..e484863709f 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -45778,6 +45778,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;
@@ -45786,6 +45788,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/iterator-with-structs.c
b/libgomp/testsuite/libgomp.c-c++-common/iterator-with-structs.c
new file mode 100644
index 00000000000..c4486f40711
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/iterator-with-structs.c
@@ -0,0 +1,24 @@
+// This test case used to ICE in the gimplifier.
+
+struct T {
+ int *p, n;
+};
+#pragma omp declare mapper(struct T t) map(t.p[ :t.n], t.n)
+
+struct S {
+ struct T t;
+};
+
+struct S s[5];
+
+int main (void)
+{
+ for (int i = 0; i < 5; i++)
+ s[i].t.n = i;
+ #pragma omp target enter data map(iterator(i=0:5), to: s[i].t)
+ for (int i = 0; i < 5; i++)
+ if (s[i].t.n != i)
+ __builtin_abort ();
+ #pragma omp target exit data map(iterator(i=0:5), release: s[i].t)
+ return 0;
+}
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 00000000000..193823744bd
--- /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 00000000000..76f00fb518a
--- /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;
+}
--
2.39.5
