https://gcc.gnu.org/g:752ef8f4f5ddf9a3372e6d5dc08df6aa05a60ea3
commit 752ef8f4f5ddf9a3372e6d5dc08df6aa05a60ea3
Author: Kwok Cheung Yeung <kcye...@baylibre.com>
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.
* 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 <a...@baylibre.com>
Diff:
---
gcc/c-family/c-omp.cc | 2 +
gcc/c/c-parser.cc | 4 +
gcc/cp/parser.cc | 4 +
gcc/cp/semantics.cc | 3 +-
.../libgomp.c-c++-common/mapper-iterators-1.c | 83 ++++++++++++++++++
.../libgomp.c-c++-common/mapper-iterators-2.c | 81 ++++++++++++++++++
.../libgomp.c-c++-common/mapper-iterators-3.c | 98 ++++++++++++++++++++++
7 files changed, 274 insertions(+), 1 deletion(-)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 7d333185fdd5..a3cd1ae0432f 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 8a5b4932dbcb..31ab7d08a2c1 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 83e02202782f..f38a32777398 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 b80bf3306775..770872ac0687 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 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;
+}
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 000000000000..9d67c38a1068
--- /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;
+}
