diff options
Diffstat (limited to 'libgomp/target.c')
| -rw-r--r-- | libgomp/target.c | 199 |
1 files changed, 192 insertions, 7 deletions
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 * |