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; +}
