This patch adds support for custom mappers with iterators (C and C++
only, as the Fortran custom mapper support has not been committed yet,
nor has support for nested mappers).
It works by propagating clause iterators onto new clauses generated by
mappers. As this occurs early in the front-end, the middle-end will
treat it the same as if the iterators were specified explicitly.
From 46a70c58505f6c0ee6af57997de81da94f0faa11 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcye...@baylibre.com>
Date: Mon, 13 Jan 2025 13:08:07 +0000
Subject: [PATCH 09/11] 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 <a...@baylibre.com>
---
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(+)
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 fe272888c51..1b68179db88 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4392,6 +4392,7 @@ omp_instantiate_mapper (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;
@@ -4476,6 +4477,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree
expr,
continue;
}
+ 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 0ecc3e88be5..1cd14b814f3 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20290,6 +20290,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;
@@ -20298,6 +20300,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 4f496433135..e819dad03ae 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -42838,6 +42838,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;
@@ -42846,6 +42848,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 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.43.0
