+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * c-omp.cc (omp_instantiate_mapper): Apply iterator to new clauses
+ generated from mapper.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* c-omp.cc (c_finish_omp_depobj): Use OMP_ITERATOR_DECL_P.
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;
}
else
{
+ OMP_CLAUSE_ITERATORS (unshared) = iterator;
*outlist = unshared;
outlist = &OMP_CLAUSE_CHAIN (unshared);
}
+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * c-parser.cc (c_parser_omp_clause_map): Apply iterator to push and
+ pop mapper clauses.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* c-parser.cc (c_parser_omp_iterators): Use macros for accessing
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;
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;
}
+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* parser.cc (cp_parser_omp_iterators): Use macros for accessing
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;
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;
}
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;
+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* testsuite/libgomp.c-c++-common/target-map-iterators-4.c: New.
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
projects / thirdparty / gcc.git / commitdiff
