diff options
Diffstat (limited to 'libgomp')
8 files changed, 549 insertions, 7 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 2689f3e..722da76 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,23 @@ +2025-08-06 Kwok Cheung Yeung <kcyeung@baylibre.com> + + * target.c (gomp_update): Call gomp_merge_iterator_maps. Free + allocated variables. + * testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New. + * testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New. + * testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New. + +2025-08-06 Kwok Cheung Yeung <kcyeung@baylibre.com> + Andrew Stubbs <ams@baylibre.com> + + * target.c (kind_to_name): New. + (gomp_merge_iterator_maps): New. + (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy + address of only the first iteration to target vars. Free allocated + variables. + * testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New. + * testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New. + * testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New. + 2025-07-21 Thomas Schwinge <tschwinge@baylibre.com> PR target/119853 diff --git a/libgomp/target.c b/libgomp/target.c index cda092b..c89c82c 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1003,6 +1003,155 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) } } +static const char * +kind_to_name (unsigned short kind, bool short_mapkind) +{ + if (short_mapkind && GOMP_MAP_IMPLICIT_P (kind)) + kind &= ~GOMP_MAP_IMPLICIT; + + switch (kind & (short_mapkind ? 0xff : 0x7)) + { + case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC"; + case GOMP_MAP_TO: return "GOMP_MAP_TO"; + case GOMP_MAP_FROM: return "GOMP_MAP_FROM"; + case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM"; + case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER"; + case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET"; + case GOMP_MAP_FORCE_PRESENT: return "GOMP_MAP_FORCE_PRESENT"; + case GOMP_MAP_DELETE: return "GOMP_MAP_DELETE"; + case GOMP_MAP_FORCE_DEVICEPTR: return "GOMP_MAP_FORCE_DEVICEPTR"; + case GOMP_MAP_DEVICE_RESIDENT: return "GOMP_MAP_DEVICE_RESIDENT"; + case GOMP_MAP_LINK: return "GOMP_MAP_LINK"; + case GOMP_MAP_IF_PRESENT: return "GOMP_MAP_IF_PRESENT"; + case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE"; + case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT"; + case GOMP_MAP_USE_DEVICE_PTR: return "GOMP_MAP_USE_DEVICE_PTR"; + case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: return "GOMP_MAP_ZERO_LEN_ARRAY_SECTION"; + case GOMP_MAP_FORCE_ALLOC: return "GOMP_MAP_FORCE_ALLOC"; + case GOMP_MAP_FORCE_TO: return "GOMP_MAP_FORCE_TO"; + case GOMP_MAP_FORCE_FROM: return "GOMP_MAP_FORCE_FROM"; + case GOMP_MAP_FORCE_TOFROM: return "GOMP_MAP_FORCE_TOFROM"; + case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT: + return "GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT"; + case GOMP_MAP_ALWAYS_TO: return "GOMP_MAP_ALWAYS_TO"; + case GOMP_MAP_ALWAYS_FROM: return "GOMP_MAP_ALWAYS_FROM"; + case GOMP_MAP_ALWAYS_TOFROM: return "GOMP_MAP_ALWAYS_TOFROM"; + case GOMP_MAP_ALWAYS_PRESENT_TO: return "GOMP_MAP_ALWAYS_PRESENT_TO"; + case GOMP_MAP_ALWAYS_PRESENT_FROM: return "GOMP_MAP_ALWAYS_PRESENT_FROM"; + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: return "GOMP_MAP_ALWAYS_PRESENT_TOFROM"; + case GOMP_MAP_STRUCT: return "GOMP_MAP_STRUCT"; + case GOMP_MAP_STRUCT_UNORD: return "GOMP_MAP_STRUCT_UNORD"; + case GOMP_MAP_ALWAYS_POINTER: return "GOMP_MAP_ALWAYS_POINTER"; + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + return "GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION"; + case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION: + return "GOMP_MAP_DELETE_ZERO_LENGTH_ARRAY_SECTION"; + case GOMP_MAP_RELEASE: return "GOMP_MAP_RELEASE"; + case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH"; + case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH"; + case GOMP_MAP_FORCE_DETACH: return "GOMP_MAP_FORCE_DETACH"; + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + return "GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION"; + default: return "unknown"; + } +} + +/* When GCC encounters a clause with an iterator, e.g.: + + #pragma omp target map (iterator(i=0:4), to: x[i]) + + it generates an array containing the number of iterations and the + address and size of each iteration. e.g.: + + void *omp_iter_data[] = { + (void *) 4, // Number of iterations + &x[0], (void *) sizeof(x[0]), + &x[1], (void *) sizeof(x[1]), + &x[2], (void *) sizeof(x[2]), + &x[3], (void *) sizeof(x[3]) + }; + + When the construct is lowered, &omp_iter_data is used as the host address + for the map (instead of &x[i]), and the size is set to SIZE_MAX to mark + the map as an iterator map. + + Map entries containing expanded iterators will be flattened and merged into + HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are + any iterators found. ITERATOR_COUNT holds the iteration count of the + iterator that generates each map (and 0 if not generated from an iterator). + HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any + merging occurs. */ + +static bool +gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, + void **kinds, size_t **iterator_count) +{ + bool iterator_p = false; + size_t map_count = 0; + unsigned short **skinds = (unsigned short **) kinds; + + for (size_t i = 0; i < *mapnum; i++) + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + map_count += iterator_array[0]; + iterator_p = true; + } + else + map_count++; + + if (!iterator_p) + return false; + + gomp_debug (1, + "Expanding iterator maps - number of map entries: %u -> %u\n", + (int) *mapnum, (int) map_count); + void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *)); + size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + unsigned short *new_kinds + = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short)); + size_t new_idx = 0; + *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + + for (size_t i = 0; i < *mapnum; i++) + { + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + size_t count = *iterator_array++; + for (size_t j = 0; j < count; j++) + { + new_hostaddrs[new_idx] = (void *) *iterator_array++; + new_sizes[new_idx] = *iterator_array++; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = j + 1; + gomp_debug (1, + "Expanding map %u <%s>: " + "hostaddrs[%u] = %p, sizes[%u] = %lu\n", + (int) i, kind_to_name (new_kinds[new_idx], true), + (int) new_idx, new_hostaddrs[new_idx], + (int) new_idx, (unsigned long) new_sizes[new_idx]); + new_idx++; + } + } + else + { + new_hostaddrs[new_idx] = (*hostaddrs)[i]; + new_sizes[new_idx] = (*sizes)[i]; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = 0; + new_idx++; + } + } + + *mapnum = map_count; + *hostaddrs = new_hostaddrs; + *sizes = new_sizes; + *kinds = new_kinds; + + return true; +} + static inline __attribute__((always_inline)) struct target_mem_desc * gomp_map_vars_internal (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, size_t mapnum, @@ -1019,6 +1168,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, const int typemask = short_mapkind ? 0xff : 0x7; struct splay_tree_s *mem_map = &devicep->mem_map; struct splay_tree_key_s cur_node; + bool iterators_p = false; + size_t *iterator_count = NULL; + if (short_mapkind) /* OpenMP */ + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; @@ -1896,14 +2050,22 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (pragma_kind & GOMP_MAP_VARS_TARGET) { + /* The target variables table is constructed with maps using iterators + unexpanded. Now that the iterator maps are expanded, we will need to + skip all expanded maps after the initial entry, otherwise subsequent + maps will be out-of-sync with their corresponding entry in the + target variables table. */ + size_t map_num = 0; for (i = 0; i < mapnum; i++) - { - cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); - gomp_copy_host2dev (devicep, aq, - (void *) (tgt->tgt_start + i * sizeof (void *)), - (void *) &cur_node.tgt_offset, sizeof (void *), - true, cbufp); - } + if (!iterator_count || iterator_count[i] <= 1) + { + cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); + gomp_copy_host2dev (devicep, aq, + (void *) (tgt->tgt_start + map_num * sizeof (void *)), + (void *) &cur_node.tgt_offset, sizeof (void *), + true, cbufp); + map_num++; + } } if (cbufp) @@ -1935,6 +2097,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } + return tgt; } @@ -2201,6 +2372,8 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t i; struct splay_tree_key_s cur_node; const int typemask = short_mapkind ? 0xff : 0x7; + bool iterators_p = false; + size_t *iterator_count = NULL; if (!devicep) return; @@ -2208,6 +2381,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, if (mapnum == 0) return; + if (short_mapkind) /* OpenMP */ + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); + gomp_mutex_lock (&devicep->lock); if (devicep->state == GOMP_DEVICE_FINALIZED) { @@ -2301,6 +2478,14 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, } } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } } static struct gomp_offload_icv_list * diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c new file mode 100644 index 0000000..b3d87f2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators. */ + +#include <stdlib.h> + +#define DIM1 8 +#define DIM2 15 + +int mkarray (int *x[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j]; + } + } + + return expected; +} + +int main (void) +{ + int *x[DIM1]; + int y; + + int expected = mkarray (x); + + #pragma omp target enter data map(to: x) + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) \ + map(from: y) + { + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + } + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c new file mode 100644 index 0000000..8569b55 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c @@ -0,0 +1,44 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays from target using map + iterators. */ + +#include <stdlib.h> + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + x[i] = (int *) malloc (DIM2 * sizeof (int)); +} + +int main (void) +{ + int *x[DIM1]; + int y, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x) + #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \ + map(from: expected) + { + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] = (i+1) * (j+1); + expected += x[i][j]; + } + } + + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c new file mode 100644 index 0000000..be30fa65d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c @@ -0,0 +1,56 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators, with multiple iterators and function calls in the iterator + expression. */ + +#include <stdlib.h> + +#define DIM1 16 +#define DIM2 15 + +int mkarrays (int *x[], int *y[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + y[i] = (int *) malloc (sizeof (int)); + *y[i] = rand (); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j] * *y[i]; + } + } + + return expected; +} + +int f (int i, int j) +{ + return i * 4 + j; +} + +int main (void) +{ + int *x[DIM1], *y[DIM1]; + int sum; + + int expected = mkarrays (x, y); + + #pragma omp target enter data map(to: x, y) + #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \ + map(iterator(i=0:DIM1), to: y[i][:1]) \ + map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j] * y[i][0]; + } + + return sum - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c new file mode 100644 index 0000000..5a4cad5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c @@ -0,0 +1,65 @@ +/* { dg-do run } */ + +/* Test target enter data and target update to the target using map + iterators. */ + +#include <stdlib.h> + +#define DIM1 8 +#define DIM2 15 + +int mkarray (int *x[]) +{ + int expected = 0; + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j]; + } + } + + return expected; +} + +int main (void) +{ + int *x[DIM1]; + int sum; + int expected = mkarray (x); + + #pragma omp target enter data map(to: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + if (sum != expected) + return 1; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] *= rand (); + expected += x[i][j]; + } + + #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + return sum != expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c new file mode 100644 index 0000000..93438d0 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c @@ -0,0 +1,58 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update from the target using map + iterators. */ + +#include <stdlib.h> + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + x[i][j] = 0; + } +} + +int main (void) +{ + int *x[DIM1]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target map(from: expected) + { + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] = (i + 1) * (j + 2); + expected += x[i][j]; + } + } + + /* Host copy of x should remain unchanged. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + if (sum != 0) + return 1; + + #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2]) + + /* Host copy should now be updated. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + return sum - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c new file mode 100644 index 0000000..a70b21c --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c @@ -0,0 +1,67 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update to the target using map + iterators with a function. */ + +#include <stdlib.h> + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + x[i][j] = rand (); + } +} + +int f (int i) +{ + return i * 2; +} + +int main (void) +{ + int *x[DIM1], x_new[DIM1][DIM2]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + + /* Update x on host. */ + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x_new[i][j] = x[i][j]; + x[i][j] = (i + 1) * (j + 2); + } + + /* Update a subset of x on target. */ + #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + /* Calculate expected value on host. */ + for (int i = 0; i < DIM1/2; i++) + for (int j = 0; j < DIM2; j++) + x_new[f (i)][j] = x[f (i)][j]; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + expected += x_new[i][j]; + + return sum - expected; +} |