From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1709) id 40CF03857431; Thu, 17 Jun 2021 13:35:19 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 40CF03857431 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Chung-Lin Tang To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-1565] libgomp: Structure element mapping for OpenMP 5.0 X-Act-Checkin: gcc X-Git-Author: Chung-Lin Tang X-Git-Refname: refs/heads/master X-Git-Oldrev: 967b46530234b4e6ad3983057705aea6c20a03c4 X-Git-Newrev: 275c736e732d29934e4d22e8f030d5aae8c12a52 Message-Id: <20210617133519.40CF03857431@sourceware.org> Date: Thu, 17 Jun 2021 13:35:19 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 17 Jun 2021 13:35:19 -0000 https://gcc.gnu.org/g:275c736e732d29934e4d22e8f030d5aae8c12a52 commit r12-1565-g275c736e732d29934e4d22e8f030d5aae8c12a52 Author: Chung-Lin Tang Date: Thu Jun 17 21:33:32 2021 +0800 libgomp: Structure element mapping for OpenMP 5.0 This patch implement OpenMP 5.0 requirements of incrementing/decrementing the reference count of a mapped structure at most once (across all elements) on a construct. This is implemented by pulling in libgomp/hashtab.h and using htab_t as a pointer set. Structure element list siblings also have pointers-to-refcounts linked together, to naturally achieve uniform increment/decrement without repeating. There are still some questions on whether using such a htab_t based set is faster/slower than using a sorted pointer array based implementation. This is to be researched on later. libgomp/ChangeLog: * hashtab.h (htab_clear): New function with initialization code factored out from... (htab_create): ...here, adjust to use htab_clear function. * libgomp.h (REFCOUNT_SPECIAL): New symbol to denote range of special refcount values, add comments. (REFCOUNT_INFINITY): Adjust definition to use REFCOUNT_SPECIAL. (REFCOUNT_LINK): Likewise. (REFCOUNT_STRUCTELEM): New special refcount range for structure element siblings. (REFCOUNT_STRUCTELEM_P): Macro for testing for structure element sibling maps. (REFCOUNT_STRUCTELEM_FLAG_FIRST): Flag to indicate first sibling. (REFCOUNT_STRUCTELEM_FLAG_LAST): Flag to indicate last sibling. (REFCOUNT_STRUCTELEM_FIRST_P): Macro to test _FIRST flag. (REFCOUNT_STRUCTELEM_LAST_P): Macro to test _LAST flag. (struct splay_tree_key_s): Add structelem_refcount and structelem_refcount_ptr fields into a union with dynamic_refcount. Add comments. (gomp_map_vars): Delete declaration. (gomp_map_vars_async): Likewise. (gomp_unmap_vars): Likewise. (gomp_unmap_vars_async): Likewise. (goacc_map_vars): New declaration. (goacc_unmap_vars): Likewise. * oacc-mem.c (acc_map_data): Adjust to use goacc_map_vars. (goacc_enter_datum): Likewise. (goacc_enter_data_internal): Likewise. * oacc-parallel.c (GOACC_parallel_keyed): Adjust to use goacc_map_vars and goacc_unmap_vars. (GOACC_data_start): Adjust to use goacc_map_vars. (GOACC_data_end): Adjust to use goacc_unmap_vars. * target.c (hash_entry_type): New typedef. (htab_alloc): New function hook for hashtab.h. (htab_free): Likewise. (htab_hash): Likewise. (htab_eq): Likewise. (hashtab.h): Add file include. (gomp_increment_refcount): New function. (gomp_decrement_refcount): Likewise. (gomp_map_vars_existing): Add refcount_set parameter, adjust to use gomp_increment_refcount. (gomp_map_fields_existing): Add refcount_set parameter, adjust calls to gomp_map_vars_existing. (gomp_map_vars_internal): Add refcount_set parameter, add local openmp_p variable to guard OpenMP specific paths, adjust calls to gomp_map_vars_existing, add structure element sibling splay_tree_key sequence creation code, adjust Fortran map case to avoid increment under OpenMP. (gomp_map_vars): Adjust to static, add refcount_set parameter, manage local refcount_set if caller passed in NULL, adjust call to gomp_map_vars_internal. (gomp_map_vars_async): Adjust and rename into... (goacc_map_vars): ...this new function, adjust call to gomp_map_vars_internal. (gomp_remove_splay_tree_key): New function with code factored out from gomp_remove_var_internal. (gomp_remove_var_internal): Add code to handle removing multiple splay_tree_key sequence for structure elements, adjust code to use gomp_remove_splay_tree_key for splay-tree key removal. (gomp_unmap_vars_internal): Add refcount_set parameter, adjust to use gomp_decrement_refcount. (gomp_unmap_vars): Adjust to static, add refcount_set parameter, manage local refcount_set if caller passed in NULL, adjust call to gomp_unmap_vars_internal. (gomp_unmap_vars_async): Adjust and rename into... (goacc_unmap_vars): ...this new function, adjust call to gomp_unmap_vars_internal. (GOMP_target): Manage refcount_set and adjust calls to gomp_map_vars and gomp_unmap_vars. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Adjust call to gomp_map_vars. (GOMP_target_data): Likewise. (GOMP_target_data_ext): Likewise. (GOMP_target_end_data): Adjust call to gomp_unmap_vars. (gomp_exit_data): Add refcount_set parameter, adjust to use gomp_decrement_refcount, adjust to queue splay-tree keys for removal after main loop. (GOMP_target_enter_exit_data): Manage refcount_set and adjust calls to gomp_map_vars and gomp_exit_data. (gomp_target_task_fn): Likewise. * testsuite/libgomp.c-c++-common/refcount-1.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-1.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-2.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-3.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-4.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-5.c: New testcase. Diff: --- libgomp/hashtab.h | 14 +- libgomp/libgomp.h | 66 ++- libgomp/oacc-mem.c | 19 +- libgomp/oacc-parallel.c | 22 +- libgomp/target.c | 442 ++++++++++++++++----- .../testsuite/libgomp.c-c++-common/refcount-1.c | 61 +++ .../testsuite/libgomp.c-c++-common/struct-elem-1.c | 29 ++ .../testsuite/libgomp.c-c++-common/struct-elem-2.c | 47 +++ .../testsuite/libgomp.c-c++-common/struct-elem-3.c | 69 ++++ .../testsuite/libgomp.c-c++-common/struct-elem-4.c | 56 +++ .../testsuite/libgomp.c-c++-common/struct-elem-5.c | 20 + 11 files changed, 709 insertions(+), 136 deletions(-) diff --git a/libgomp/hashtab.h b/libgomp/hashtab.h index 25d6d94c571..26d80813436 100644 --- a/libgomp/hashtab.h +++ b/libgomp/hashtab.h @@ -224,6 +224,15 @@ htab_mod_m2 (hashval_t hash, htab_t htab) return 1 + htab_mod_1 (hash, p->prime - 2, p->inv_m2, p->shift); } +static inline htab_t +htab_clear (htab_t htab) +{ + htab->n_elements = 0; + htab->n_deleted = 0; + memset (htab->entries, 0, htab->size * sizeof (hash_entry_type)); + return htab; +} + /* Create hash table of size SIZE. */ static htab_t @@ -238,11 +247,8 @@ htab_create (size_t size) result = (htab_t) htab_alloc (sizeof (struct htab) + size * sizeof (hash_entry_type)); result->size = size; - result->n_elements = 0; - result->n_deleted = 0; result->size_prime_index = size_prime_index; - memset (result->entries, 0, size * sizeof (hash_entry_type)); - return result; + return htab_clear (result); } /* Similar to htab_find_slot, but without several unwanted side effects: diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index ef1bb4907b6..8d25dc8e2a8 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1012,11 +1012,35 @@ struct target_mem_desc { struct target_var_desc list[]; }; +/* Special value for refcount - mask to indicate existence of special + values. Right now we allocate 3 bits. */ +#define REFCOUNT_SPECIAL (~(uintptr_t) 0x7) + /* Special value for refcount - infinity. */ -#define REFCOUNT_INFINITY (~(uintptr_t) 0) +#define REFCOUNT_INFINITY (REFCOUNT_SPECIAL | 0) /* Special value for refcount - tgt_offset contains target address of the artificial pointer to "omp declare target link" object. */ -#define REFCOUNT_LINK (~(uintptr_t) 1) +#define REFCOUNT_LINK (REFCOUNT_SPECIAL | 1) + +/* Special value for refcount - structure element sibling list items. + All such key refounts have REFCOUNT_STRUCTELEM bits set, with _FLAG_FIRST + and _FLAG_LAST indicating first and last in the created sibling sequence. */ +#define REFCOUNT_STRUCTELEM (REFCOUNT_SPECIAL | 4) +#define REFCOUNT_STRUCTELEM_P(V) \ + (((V) & REFCOUNT_STRUCTELEM) == REFCOUNT_STRUCTELEM) +/* The first leading key with _FLAG_FIRST set houses the actual reference count + in the structelem_refcount field. Other siblings point to this counter value + through its structelem_refcount_ptr field. */ +#define REFCOUNT_STRUCTELEM_FLAG_FIRST (1) +/* The last key in the sibling sequence has this set. This is required to + indicate the sequence boundary, when we remove the structure sibling list + from the map. */ +#define REFCOUNT_STRUCTELEM_FLAG_LAST (2) + +#define REFCOUNT_STRUCTELEM_FIRST_P(V) \ + (REFCOUNT_STRUCTELEM_P (V) && ((V) & REFCOUNT_STRUCTELEM_FLAG_FIRST)) +#define REFCOUNT_STRUCTELEM_LAST_P(V) \ + (REFCOUNT_STRUCTELEM_P (V) && ((V) & REFCOUNT_STRUCTELEM_FLAG_LAST)) /* Special offset values. */ #define OFFSET_INLINED (~(uintptr_t) 0) @@ -1044,8 +1068,22 @@ struct splay_tree_key_s { uintptr_t tgt_offset; /* Reference count. */ uintptr_t refcount; - /* Dynamic reference count. */ - uintptr_t dynamic_refcount; + union { + /* Dynamic reference count. */ + uintptr_t dynamic_refcount; + + /* Unified reference count for structure element siblings, this is used + when REFCOUNT_STRUCTELEM_FIRST_P(k->refcount) == true, the first sibling + in a structure element sibling list item sequence. */ + uintptr_t structelem_refcount; + + /* When REFCOUNT_STRUCTELEM_P (k->refcount) == true, this field points + into the (above) structelem_refcount field of the _FIRST splay_tree_key, + the first key in the created sequence. All structure element siblings + share a single refcount in this manner. Since these two fields won't be + used at the same time, they are stashed in a union. */ + uintptr_t *structelem_refcount_ptr; + }; struct splay_tree_aux *aux; }; @@ -1200,19 +1238,13 @@ extern void gomp_attach_pointer (struct gomp_device_descr *, extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); - -extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *, - size_t, void **, void **, - size_t *, void *, bool, - enum gomp_map_vars_kind); -extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *, - struct goacc_asyncqueue *, - size_t, void **, void **, - size_t *, void *, bool, - enum gomp_map_vars_kind); -extern void gomp_unmap_vars (struct target_mem_desc *, bool); -extern void gomp_unmap_vars_async (struct target_mem_desc *, bool, - struct goacc_asyncqueue *); +extern struct target_mem_desc *goacc_map_vars (struct gomp_device_descr *, + struct goacc_asyncqueue *, + size_t, void **, void **, + size_t *, void *, bool, + enum gomp_map_vars_kind); +extern void goacc_unmap_vars (struct target_mem_desc *, bool, + struct goacc_asyncqueue *); extern void gomp_init_device (struct gomp_device_descr *); extern bool gomp_fini_device (struct gomp_device_descr *); extern void gomp_unload_device (struct gomp_device_descr *); diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 123fe154089..c21508f3739 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -402,9 +402,8 @@ acc_map_data (void *h, void *d, size_t s) gomp_mutex_unlock (&acc_dev->lock); struct target_mem_desc *tgt - = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, - &kinds, true, - GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA); + = goacc_map_vars (acc_dev, NULL, mapnum, &hostaddrs, &devaddrs, &sizes, + &kinds, true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); assert (tgt->list_count == 1); splay_tree_key n = tgt->list[0].key; @@ -572,9 +571,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) goacc_aq aq = get_goacc_asyncqueue (async); struct target_mem_desc *tgt - = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, - kinds, true, (GOMP_MAP_VARS_OPENACC - | GOMP_MAP_VARS_ENTER_DATA)); + = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, + kinds, true, GOMP_MAP_VARS_ENTER_DATA); assert (tgt); assert (tgt->list_count == 1); n = tgt->list[0].key; @@ -1070,7 +1068,7 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) } /* Map variables for OpenACC "enter data". We can't just call - gomp_map_vars_async once, because individual mapped variables might have + goacc_map_vars once, because individual mapped variables might have "exit data" called for them at different times. */ static void @@ -1202,10 +1200,9 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); struct target_mem_desc *tgt - = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, - &sizes[i], &kinds[i], true, - (GOMP_MAP_VARS_OPENACC - | GOMP_MAP_VARS_ENTER_DATA)); + = goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL, + &sizes[i], &kinds[i], true, + GOMP_MAP_VARS_ENTER_DATA); assert (tgt); gomp_mutex_lock (&acc_dev->lock); diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 939e09f2b0c..83625ba8a8e 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -290,8 +290,8 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), goacc_aq aq = get_goacc_asyncqueue (async); - tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, - true, GOMP_MAP_VARS_OPENACC); + tgt = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, + true, 0); if (profiling_p) { prof_info.event_type = acc_ev_enter_data_end; @@ -300,7 +300,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), goacc_profiling_dispatch (&prof_info, &enter_exit_data_event_info, &api_info); } - + devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, i); @@ -321,11 +321,8 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), &api_info); } - /* If running synchronously, unmap immediately. */ - if (aq == NULL) - gomp_unmap_vars (tgt, true); - else - gomp_unmap_vars_async (tgt, true, aq); + /* If running synchronously (aq == NULL), this will unmap immediately. */ + goacc_unmap_vars (tgt, true, aq); if (profiling_p) { @@ -456,8 +453,7 @@ GOACC_data_start (int flags_m, size_t mapnum, { prof_info.device_type = acc_device_host; api_info.device_type = prof_info.device_type; - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); + tgt = goacc_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, 0); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -465,8 +461,8 @@ GOACC_data_start (int flags_m, size_t mapnum, } gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__); - tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, - GOMP_MAP_VARS_OPENACC); + tgt = goacc_map_vars (acc_dev, NULL, mapnum, hostaddrs, NULL, sizes, kinds, + true, 0); gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -542,7 +538,7 @@ GOACC_data_end (void) gomp_debug (0, " %s: restore mappings\n", __FUNCTION__); thr->mapped_data = tgt->prev; - gomp_unmap_vars (tgt, true); + goacc_unmap_vars (tgt, true, NULL); gomp_debug (0, " %s: mappings restored\n", __FUNCTION__); if (profiling_p) diff --git a/libgomp/target.c b/libgomp/target.c index 2150e5d79b2..bb09d501dd6 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -44,6 +44,23 @@ #include "plugin-suffix.h" #endif +typedef uintptr_t *hash_entry_type; +static inline void * htab_alloc (size_t size) { return gomp_malloc (size); } +static inline void htab_free (void *ptr) { free (ptr); } +#include "hashtab.h" + +static inline hashval_t +htab_hash (hash_entry_type element) +{ + return hash_pointer ((void *) element); +} + +static inline bool +htab_eq (hash_entry_type x, hash_entry_type y) +{ + return x == y; +} + #define FIELD_TGT_EMPTY (~(size_t) 0) static void gomp_target_init (void); @@ -360,6 +377,113 @@ gomp_free_device_memory (struct gomp_device_descr *devicep, void *devptr) } } +/* Increment reference count of a splay_tree_key region K by 1. + If REFCOUNT_SET != NULL, use it to track already seen refcounts, and only + increment the value if refcount is not yet contained in the set (used for + OpenMP 5.0, which specifies that a region's refcount is adjusted at most + once for each construct). */ + +static inline void +gomp_increment_refcount (splay_tree_key k, htab_t *refcount_set) +{ + if (k == NULL || k->refcount == REFCOUNT_INFINITY) + return; + + uintptr_t *refcount_ptr = &k->refcount; + + if (REFCOUNT_STRUCTELEM_FIRST_P (k->refcount)) + refcount_ptr = &k->structelem_refcount; + else if (REFCOUNT_STRUCTELEM_P (k->refcount)) + refcount_ptr = k->structelem_refcount_ptr; + + if (refcount_set) + { + if (htab_find (*refcount_set, refcount_ptr)) + return; + uintptr_t **slot = htab_find_slot (refcount_set, refcount_ptr, INSERT); + *slot = refcount_ptr; + } + + *refcount_ptr += 1; + return; +} + +/* Decrement reference count of a splay_tree_key region K by 1, or if DELETE_P + is true, set reference count to zero. If REFCOUNT_SET != NULL, use it to + track already seen refcounts, and only adjust the value if refcount is not + yet contained in the set (like gomp_increment_refcount). + + Return out-values: set *DO_COPY to true if we set the refcount to zero, or + it is already zero and we know we decremented it earlier. This signals that + associated maps should be copied back to host. + + *DO_REMOVE is set to true when we this is the first handling of this refcount + and we are setting it to zero. This signals a removal of this key from the + splay-tree map. + + Copy and removal are separated due to cases like handling of structure + elements, e.g. each map of a structure element representing a possible copy + out of a structure field has to be handled individually, but we only signal + removal for one (the first encountered) sibing map. */ + +static inline void +gomp_decrement_refcount (splay_tree_key k, htab_t *refcount_set, bool delete_p, + bool *do_copy, bool *do_remove) +{ + if (k == NULL || k->refcount == REFCOUNT_INFINITY) + { + *do_copy = *do_remove = false; + return; + } + + uintptr_t *refcount_ptr = &k->refcount; + + if (REFCOUNT_STRUCTELEM_FIRST_P (k->refcount)) + refcount_ptr = &k->structelem_refcount; + else if (REFCOUNT_STRUCTELEM_P (k->refcount)) + refcount_ptr = k->structelem_refcount_ptr; + + bool new_encountered_refcount; + bool set_to_zero = false; + bool is_zero = false; + + uintptr_t orig_refcount = *refcount_ptr; + + if (refcount_set) + { + if (htab_find (*refcount_set, refcount_ptr)) + { + new_encountered_refcount = false; + goto end; + } + + uintptr_t **slot = htab_find_slot (refcount_set, refcount_ptr, INSERT); + *slot = refcount_ptr; + new_encountered_refcount = true; + } + else + /* If no refcount_set being used, assume all keys are being decremented + for the first time. */ + new_encountered_refcount = true; + + if (delete_p) + *refcount_ptr = 0; + else if (*refcount_ptr > 0) + *refcount_ptr -= 1; + + end: + if (*refcount_ptr == 0) + { + if (orig_refcount > 0) + set_to_zero = true; + + is_zero = true; + } + + *do_copy = (set_to_zero || (!new_encountered_refcount && is_zero)); + *do_remove = (new_encountered_refcount && set_to_zero); +} + /* Handle the case where gomp_map_lookup, splay_tree_lookup or gomp_map_0len_lookup found oldn for newn. Helper function of gomp_map_vars. */ @@ -369,7 +493,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, splay_tree_key oldn, splay_tree_key newn, struct target_var_desc *tgt_var, unsigned char kind, bool always_to_flag, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, + htab_t *refcount_set) { assert (kind != GOMP_MAP_ATTACH); @@ -398,8 +523,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, (void *) newn->host_start, newn->host_end - newn->host_start, cbuf); - if (oldn->refcount != REFCOUNT_INFINITY) - oldn->refcount++; + gomp_increment_refcount (oldn, refcount_set); } static int @@ -453,7 +577,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, splay_tree_key n, size_t first, size_t i, void **hostaddrs, size_t *sizes, void *kinds, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -471,7 +595,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) { gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], - kind & typemask, false, cbuf); + kind & typemask, false, cbuf, refcount_set); return; } if (sizes[i] == 0) @@ -487,7 +611,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, == n2->tgt_offset - n->tgt_offset) { gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], - kind & typemask, false, cbuf); + kind & typemask, false, cbuf, refcount_set); return; } } @@ -499,7 +623,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) { gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], - kind & typemask, false, cbuf); + kind & typemask, false, cbuf, refcount_set); return; } } @@ -671,11 +795,13 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, size_t mapnum, void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, bool short_mapkind, + htab_t *refcount_set, enum gomp_map_vars_kind pragma_kind) { size_t i, tgt_align, tgt_size, not_found_cnt = 0; bool has_firstprivate = false; bool has_always_ptrset = false; + bool openmp_p = (pragma_kind & GOMP_MAP_VARS_OPENACC) == 0; const int rshift = short_mapkind ? 8 : 3; const int typemask = short_mapkind ? 0xff : 0x7; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -813,7 +939,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } for (i = first; i <= last; i++) gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs, - sizes, kinds, NULL); + sizes, kinds, NULL, refcount_set); i--; continue; } @@ -909,7 +1035,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } } gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i], - kind & typemask, always_to_cnt > 0, NULL); + kind & typemask, always_to_cnt > 0, NULL, + refcount_set); i += always_to_cnt; } else @@ -1022,6 +1149,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, splay_tree_node array = tgt->array; size_t j, field_tgt_offset = 0, field_tgt_clear = FIELD_TGT_EMPTY; uintptr_t field_tgt_base = 0; + splay_tree_key field_tgt_structelem_first = NULL; for (i = 0; i < mapnum; i++) if (has_always_ptrset @@ -1064,8 +1192,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[j].copy_from = false; tgt->list[j].always_copy_from = false; tgt->list[j].is_attach = false; - if (k->refcount != REFCOUNT_INFINITY) - k->refcount++; + gomp_increment_refcount (k, refcount_set); gomp_map_pointer (k->tgt, aq, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset + ((uintptr_t) hostaddrs[j] @@ -1153,13 +1280,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, field_tgt_base = (uintptr_t) hostaddrs[first]; field_tgt_offset = tgt_size; field_tgt_clear = last; + field_tgt_structelem_first = NULL; tgt_size += cur_node.host_end - (uintptr_t) hostaddrs[first]; continue; } for (i = first; i <= last; i++) gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs, - sizes, kinds, cbufp); + sizes, kinds, cbufp, refcount_set); i--; continue; case GOMP_MAP_ALWAYS_POINTER: @@ -1236,7 +1364,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, splay_tree_key n = splay_tree_lookup (mem_map, k); if (n && n->refcount != REFCOUNT_LINK) gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], - kind & typemask, false, cbufp); + kind & typemask, false, cbufp, + refcount_set); else { k->aux = NULL; @@ -1252,10 +1381,33 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, size_t align = (size_t) 1 << (kind >> rshift); tgt->list[i].key = k; k->tgt = tgt; + k->refcount = 0; + k->dynamic_refcount = 0; if (field_tgt_clear != FIELD_TGT_EMPTY) { k->tgt_offset = k->host_start - field_tgt_base + field_tgt_offset; + if (openmp_p) + { + k->refcount = REFCOUNT_STRUCTELEM; + if (field_tgt_structelem_first == NULL) + { + /* Set to first structure element of sequence. */ + k->refcount |= REFCOUNT_STRUCTELEM_FLAG_FIRST; + field_tgt_structelem_first = k; + } + else + /* Point to refcount of leading element, but do not + increment again. */ + k->structelem_refcount_ptr + = &field_tgt_structelem_first->structelem_refcount; + + if (i == field_tgt_clear) + { + k->refcount |= REFCOUNT_STRUCTELEM_FLAG_LAST; + field_tgt_structelem_first = NULL; + } + } if (i == field_tgt_clear) field_tgt_clear = FIELD_TGT_EMPTY; } @@ -1265,14 +1417,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset = tgt_size; tgt_size += k->host_end - k->host_start; } + /* First increment, from 0 to 1. gomp_increment_refcount + encapsulates the different increment cases, so use this + instead of directly setting 1 during initialization. */ + gomp_increment_refcount (k, refcount_set); + tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); tgt->list[i].always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind & typemask); tgt->list[i].is_attach = false; tgt->list[i].offset = 0; tgt->list[i].length = k->host_end - k->host_start; - k->refcount = 1; - k->dynamic_refcount = 0; tgt->refcount++; array->left = NULL; array->right = NULL; @@ -1328,8 +1483,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[j].always_copy_from = false; tgt->list[j].is_attach = false; tgt->list[i].has_null_ptr_assoc |= !(*(void **) hostaddrs[j]); - if (k->refcount != REFCOUNT_INFINITY) - k->refcount++; + /* For OpenMP, the use of refcount_sets causes + errors if we set k->refcount = 1 above but also + increment it again here, for decrementing will + not properly match, since we decrement only once + for each key's refcount. Therefore avoid this + increment for OpenMP constructs. */ + if (!openmp_p) + gomp_increment_refcount (k, refcount_set); gomp_map_pointer (tgt, aq, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset @@ -1426,24 +1587,41 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, return tgt; } -attribute_hidden struct target_mem_desc * +static struct target_mem_desc * gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, - bool short_mapkind, enum gomp_map_vars_kind pragma_kind) + bool short_mapkind, htab_t *refcount_set, + enum gomp_map_vars_kind pragma_kind) { - return gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, pragma_kind); + /* This management of a local refcount_set is for convenience of callers + who do not share a refcount_set over multiple map/unmap uses. */ + htab_t local_refcount_set = NULL; + if (refcount_set == NULL) + { + local_refcount_set = htab_create (mapnum); + refcount_set = &local_refcount_set; + } + + struct target_mem_desc *tgt; + tgt = gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs, + sizes, kinds, short_mapkind, refcount_set, + pragma_kind); + if (local_refcount_set) + htab_free (local_refcount_set); + + return tgt; } attribute_hidden struct target_mem_desc * -gomp_map_vars_async (struct gomp_device_descr *devicep, - struct goacc_asyncqueue *aq, size_t mapnum, - void **hostaddrs, void **devaddrs, size_t *sizes, - void *kinds, bool short_mapkind, - enum gomp_map_vars_kind pragma_kind) +goacc_map_vars (struct gomp_device_descr *devicep, + struct goacc_asyncqueue *aq, size_t mapnum, + void **hostaddrs, void **devaddrs, size_t *sizes, + void *kinds, bool short_mapkind, + enum gomp_map_vars_kind pragma_kind) { return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, pragma_kind); + sizes, kinds, short_mapkind, NULL, + GOMP_MAP_VARS_OPENACC | pragma_kind); } static void @@ -1481,22 +1659,56 @@ gomp_unref_tgt_void (void *ptr) (void) gomp_unref_tgt (ptr); } -static inline __attribute__((always_inline)) bool -gomp_remove_var_internal (struct gomp_device_descr *devicep, splay_tree_key k, - struct goacc_asyncqueue *aq) +static void +gomp_remove_splay_tree_key (splay_tree sp, splay_tree_key k) { - bool is_tgt_unmapped = false; - splay_tree_remove (&devicep->mem_map, k); + splay_tree_remove (sp, k); if (k->aux) { if (k->aux->link_key) - splay_tree_insert (&devicep->mem_map, - (splay_tree_node) k->aux->link_key); + splay_tree_insert (sp, (splay_tree_node) k->aux->link_key); if (k->aux->attach_count) free (k->aux->attach_count); free (k->aux); k->aux = NULL; } +} + +static inline __attribute__((always_inline)) bool +gomp_remove_var_internal (struct gomp_device_descr *devicep, splay_tree_key k, + struct goacc_asyncqueue *aq) +{ + bool is_tgt_unmapped = false; + + if (REFCOUNT_STRUCTELEM_P (k->refcount)) + { + if (REFCOUNT_STRUCTELEM_FIRST_P (k->refcount) == false) + /* Infer the splay_tree_key of the first structelem key using the + pointer to the first structleme_refcount. */ + k = (splay_tree_key) ((char *) k->structelem_refcount_ptr + - offsetof (struct splay_tree_key_s, + structelem_refcount)); + assert (REFCOUNT_STRUCTELEM_FIRST_P (k->refcount)); + + /* The array created by gomp_map_vars is an array of splay_tree_nodes, + with the splay_tree_keys embedded inside. */ + splay_tree_node node = + (splay_tree_node) ((char *) k + - offsetof (struct splay_tree_node_s, key)); + while (true) + { + /* Starting from the _FIRST key, and continue for all following + sibling keys. */ + gomp_remove_splay_tree_key (&devicep->mem_map, k); + if (REFCOUNT_STRUCTELEM_LAST_P (k->refcount)) + break; + else + k = &(++node)->key; + } + } + else + gomp_remove_splay_tree_key (&devicep->mem_map, k); + if (aq) devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void, (void *) k->tgt); @@ -1530,7 +1742,7 @@ gomp_remove_var_async (struct gomp_device_descr *devicep, splay_tree_key k, static inline __attribute__((always_inline)) void gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, - struct goacc_asyncqueue *aq) + htab_t *refcount_set, struct goacc_asyncqueue *aq) { struct gomp_device_descr *devicep = tgt->device_descr; @@ -1573,23 +1785,17 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, if (tgt->list[i].is_attach) continue; - bool do_unmap = false; - if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) - k->refcount--; - else if (k->refcount == 1) - { - k->refcount--; - do_unmap = true; - } + bool do_copy, do_remove; + gomp_decrement_refcount (k, refcount_set, false, &do_copy, &do_remove); - if ((do_unmap && do_copyfrom && tgt->list[i].copy_from) + if ((do_copy && do_copyfrom && tgt->list[i].copy_from) || tgt->list[i].always_copy_from) gomp_copy_dev2host (devicep, aq, (void *) (k->host_start + tgt->list[i].offset), (void *) (k->tgt->tgt_start + k->tgt_offset + tgt->list[i].offset), tgt->list[i].length); - if (do_unmap) + if (do_remove) { struct target_mem_desc *k_tgt = k->tgt; bool is_tgt_unmapped = gomp_remove_var (devicep, k); @@ -1610,17 +1816,30 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, gomp_mutex_unlock (&devicep->lock); } -attribute_hidden void -gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom) +static void +gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, + htab_t *refcount_set) { - gomp_unmap_vars_internal (tgt, do_copyfrom, NULL); + /* This management of a local refcount_set is for convenience of callers + who do not share a refcount_set over multiple map/unmap uses. */ + htab_t local_refcount_set = NULL; + if (refcount_set == NULL) + { + local_refcount_set = htab_create (tgt->list_count); + refcount_set = &local_refcount_set; + } + + gomp_unmap_vars_internal (tgt, do_copyfrom, refcount_set, NULL); + + if (local_refcount_set) + htab_free (local_refcount_set); } attribute_hidden void -gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom, - struct goacc_asyncqueue *aq) +goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, + struct goacc_asyncqueue *aq) { - gomp_unmap_vars_internal (tgt, do_copyfrom, aq); + gomp_unmap_vars_internal (tgt, do_copyfrom, NULL, aq); } static void @@ -2130,12 +2349,15 @@ GOMP_target (int device, void (*fn) (void *), const void *unused, || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))) return gomp_target_fallback (fn, hostaddrs, devicep); + htab_t refcount_set = htab_create (mapnum); struct target_mem_desc *tgt_vars = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, - GOMP_MAP_VARS_TARGET); + &refcount_set, GOMP_MAP_VARS_TARGET); devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start, NULL); - gomp_unmap_vars (tgt_vars, true); + htab_clear (refcount_set); + gomp_unmap_vars (tgt_vars, true, &refcount_set); + htab_free (refcount_set); } static inline unsigned int @@ -2269,6 +2491,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, } struct target_mem_desc *tgt_vars; + htab_t refcount_set = NULL; + if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) { if (!fpc_done) @@ -2285,13 +2509,21 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, tgt_vars = NULL; } else - tgt_vars = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, - true, GOMP_MAP_VARS_TARGET); + { + refcount_set = htab_create (mapnum); + tgt_vars = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, + true, &refcount_set, GOMP_MAP_VARS_TARGET); + } devicep->run_func (devicep->target_id, fn_addr, tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs, args); if (tgt_vars) - gomp_unmap_vars (tgt_vars, true); + { + htab_clear (refcount_set); + gomp_unmap_vars (tgt_vars, true, &refcount_set); + } + if (refcount_set) + htab_free (refcount_set); } /* Host fallback for GOMP_target_data{,_ext} routines. */ @@ -2314,7 +2546,7 @@ gomp_target_data_fallback (struct gomp_device_descr *devicep) would get out of sync. */ struct target_mem_desc *tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, - GOMP_MAP_VARS_DATA); + NULL, GOMP_MAP_VARS_DATA); tgt->prev = icv->target_data; icv->target_data = tgt; } @@ -2333,7 +2565,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum, struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, - GOMP_MAP_VARS_DATA); + NULL, GOMP_MAP_VARS_DATA); struct gomp_task_icv *icv = gomp_icv (true); tgt->prev = icv->target_data; icv->target_data = tgt; @@ -2352,7 +2584,7 @@ GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs, struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true, - GOMP_MAP_VARS_DATA); + NULL, GOMP_MAP_VARS_DATA); struct gomp_task_icv *icv = gomp_icv (true); tgt->prev = icv->target_data; icv->target_data = tgt; @@ -2366,7 +2598,7 @@ GOMP_target_end_data (void) { struct target_mem_desc *tgt = icv->target_data; icv->target_data = tgt->prev; - gomp_unmap_vars (tgt, true); + gomp_unmap_vars (tgt, true, NULL); } } @@ -2465,7 +2697,8 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs, static void gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, - void **hostaddrs, size_t *sizes, unsigned short *kinds) + void **hostaddrs, size_t *sizes, unsigned short *kinds, + htab_t *refcount_set) { const int typemask = 0xff; size_t i; @@ -2489,6 +2722,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, false, NULL); } + int nrmvars = 0; + splay_tree_key remove_vars[mapnum]; + for (i = 0; i < mapnum; i++) { struct splay_tree_key_s cur_node; @@ -2510,22 +2746,32 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, if (!k) continue; - if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY) - k->refcount--; - if ((kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION) - && k->refcount != REFCOUNT_INFINITY) - k->refcount = 0; + bool delete_p = (kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION); + bool do_copy, do_remove; + gomp_decrement_refcount (k, refcount_set, delete_p, &do_copy, + &do_remove); - if ((kind == GOMP_MAP_FROM && k->refcount == 0) + if ((kind == GOMP_MAP_FROM && do_copy) || kind == GOMP_MAP_ALWAYS_FROM) gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start, (void *) (k->tgt->tgt_start + k->tgt_offset + cur_node.host_start - k->host_start), cur_node.host_end - cur_node.host_start); - if (k->refcount == 0) - gomp_remove_var (devicep, k); + + /* Structure elements lists are removed altogether at once, which + may cause immediate deallocation of the target_mem_desc, causing + errors if we still have following element siblings to copy back. + While we're at it, it also seems more disciplined to simply + queue all removals together for processing below. + + Structured block unmapping (i.e. gomp_unmap_vars_internal) should + not have this problem, since they maintain an additional + tgt->refcount = 1 reference to the target_mem_desc to start with. + */ + if (do_remove) + remove_vars[nrmvars++] = k; break; case GOMP_MAP_DETACH: @@ -2537,6 +2783,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, } } + for (int i = 0; i < nrmvars; i++) + gomp_remove_var (devicep, remove_vars[i]); + gomp_mutex_unlock (&devicep->lock); } @@ -2616,6 +2865,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, } } + htab_t refcount_set = htab_create (mapnum); + /* The variables are mapped separately such that they can be released independently. */ size_t i, j; @@ -2624,7 +2875,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT) { gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i], - &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); + &kinds[i], true, &refcount_set, + GOMP_MAP_VARS_ENTER_DATA); i += sizes[i]; } else if ((kinds[i] & 0xff) == GOMP_MAP_TO_PSET) @@ -2634,7 +2886,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, && !GOMP_MAP_ALWAYS_POINTER_P (get_kind (true, kinds, j) & 0xff)) break; gomp_map_vars (devicep, j-i, &hostaddrs[i], NULL, &sizes[i], - &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); + &kinds[i], true, &refcount_set, + GOMP_MAP_VARS_ENTER_DATA); i += j - i - 1; } else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH) @@ -2642,14 +2895,15 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, /* An attach operation must be processed together with the mapped base-pointer list item. */ gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i], - true, GOMP_MAP_VARS_ENTER_DATA); + true, &refcount_set, GOMP_MAP_VARS_ENTER_DATA); i += 1; } else gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i], - true, GOMP_MAP_VARS_ENTER_DATA); + true, &refcount_set, GOMP_MAP_VARS_ENTER_DATA); else - gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds); + gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds, &refcount_set); + htab_free (refcount_set); } bool @@ -2674,7 +2928,7 @@ gomp_target_task_fn (void *data) if (ttask->state == GOMP_TARGET_TASK_FINISHED) { if (ttask->tgt) - gomp_unmap_vars (ttask->tgt, true); + gomp_unmap_vars (ttask->tgt, true, NULL); return false; } @@ -2688,7 +2942,7 @@ gomp_target_task_fn (void *data) { ttask->tgt = gomp_map_vars (devicep, ttask->mapnum, ttask->hostaddrs, NULL, ttask->sizes, ttask->kinds, true, - GOMP_MAP_VARS_TARGET); + NULL, GOMP_MAP_VARS_TARGET); actual_arguments = (void *) ttask->tgt->tgt_start; } ttask->state = GOMP_TARGET_TASK_READY_TO_RUN; @@ -2707,21 +2961,27 @@ gomp_target_task_fn (void *data) if (ttask->flags & GOMP_TARGET_FLAG_UPDATE) gomp_update (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes, ttask->kinds, true); - else if ((ttask->flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0) - for (i = 0; i < ttask->mapnum; i++) - if ((ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT) - { - gomp_map_vars (devicep, ttask->sizes[i] + 1, &ttask->hostaddrs[i], - NULL, &ttask->sizes[i], &ttask->kinds[i], true, - GOMP_MAP_VARS_ENTER_DATA); - i += ttask->sizes[i]; - } - else - gomp_map_vars (devicep, 1, &ttask->hostaddrs[i], NULL, &ttask->sizes[i], - &ttask->kinds[i], true, GOMP_MAP_VARS_ENTER_DATA); else - gomp_exit_data (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes, - ttask->kinds); + { + htab_t refcount_set = htab_create (ttask->mapnum); + if ((ttask->flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0) + for (i = 0; i < ttask->mapnum; i++) + if ((ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT) + { + gomp_map_vars (devicep, ttask->sizes[i] + 1, &ttask->hostaddrs[i], + NULL, &ttask->sizes[i], &ttask->kinds[i], true, + &refcount_set, GOMP_MAP_VARS_ENTER_DATA); + i += ttask->sizes[i]; + } + else + gomp_map_vars (devicep, 1, &ttask->hostaddrs[i], NULL, &ttask->sizes[i], + &ttask->kinds[i], true, &refcount_set, + GOMP_MAP_VARS_ENTER_DATA); + else + gomp_exit_data (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes, + ttask->kinds, &refcount_set); + htab_free (refcount_set); + } return false; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c b/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c new file mode 100644 index 00000000000..5ccd908aa85 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c @@ -0,0 +1,61 @@ +#include +#include + +int main (void) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + unsigned int a = 0xcdcdcdcd; + #pragma omp target enter data map (to:a) + + a = 0xabababab; + unsigned char *p = (unsigned char *) &a; + unsigned char *q = p + 2; + + #pragma omp target enter data map (alloc:p[:1], q[:1]) + + if (d != id) + { + if (!omp_target_is_present (&a, d)) + abort (); + if (!omp_target_is_present (&p[0], d)) + abort (); + if (!omp_target_is_present (&q[0], d)) + abort (); + } + + #pragma omp target exit data map (release:a) + + if (d != id) + { + if (!omp_target_is_present (&a, d)) + abort (); + if (!omp_target_is_present (&p[0], d)) + abort (); + if (!omp_target_is_present (&q[0], d)) + abort (); + } + + #pragma omp target exit data map (from:q[:1]) + + if (d != id) + { + if (omp_target_is_present (&a, d)) + abort (); + if (omp_target_is_present (&p[0], d)) + abort (); + if (omp_target_is_present (&q[0], d)) + abort (); + + if (q[0] != 0xcd) + abort (); + if (p[0] != 0xab) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c new file mode 100644 index 00000000000..5f40fd7830c --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c @@ -0,0 +1,29 @@ +#include +#include + +struct S +{ + int a, b; +}; +typedef struct S S; + +int main (void) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + S s; + #pragma omp target enter data map (alloc: s.a, s.b) + #pragma omp target exit data map (release: s.b) + + /* OpenMP 5.0 structure element mapping rules describe that elements of same + structure variable should allocate/deallocate in a uniform fashion, so + "s.a" should be removed together by above 'exit data'. */ + if (d != id && omp_target_is_present (&s.a, d)) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c new file mode 100644 index 00000000000..c50b299c607 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c @@ -0,0 +1,47 @@ +#include +#include + +struct S +{ + int a, b, c, d; +}; +typedef struct S S; + +int main (void) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + S s; + #pragma omp target enter data map (alloc: s.a, s.b, s.c, s.d) + #pragma omp target enter data map (alloc: s.c) + #pragma omp target enter data map (alloc: s.b, s.d) + #pragma omp target enter data map (alloc: s.a, s.c, s.b) + + #pragma omp target exit data map (release: s.a) + #pragma omp target exit data map (release: s.d) + #pragma omp target exit data map (release: s.c) + #pragma omp target exit data map (release: s.b) + + /* OpenMP 5.0 structure element mapping rules describe that elements of same + structure variable should allocate/deallocate in a uniform fashion, so + all elements of 's' should be removed together by above 'exit data's. */ + if (d != id) + { + if (omp_target_is_present (&s, d)) + abort (); + if (omp_target_is_present (&s.a, d)) + abort (); + if (omp_target_is_present (&s.b, d)) + abort (); + if (omp_target_is_present (&s.c, d)) + abort (); + if (omp_target_is_present (&s.d, d)) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c new file mode 100644 index 00000000000..e2b6a6a2b28 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c @@ -0,0 +1,69 @@ +#include +#include + +struct S +{ + int a, b, c, d; +}; +typedef struct S S; + +int main (void) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + S s; + + #pragma omp target enter data map (alloc: s) + #pragma omp target enter data map (alloc: s) + + #pragma omp target exit data map (release: s.a) + #pragma omp target exit data map (release: s.b) + + /* OpenMP 5.0 structure element mapping rules describe that elements of same + structure variable should allocate/deallocate in a uniform fashion, so + all elements of 's' should be removed together by above 'exit data's. */ + if (d != id) + { + if (omp_target_is_present (&s, d)) + abort (); + if (omp_target_is_present (&s.a, d)) + abort (); + if (omp_target_is_present (&s.b, d)) + abort (); + if (omp_target_is_present (&s.c, d)) + abort (); + if (omp_target_is_present (&s.d, d)) + abort (); + } + + #pragma omp target enter data map (alloc: s.a, s.b) + #pragma omp target enter data map (alloc: s.a) + #pragma omp target enter data map (alloc: s.b) + + #pragma omp target exit data map (release: s) + #pragma omp target exit data map (release: s) + #pragma omp target exit data map (release: s) + + /* OpenMP 5.0 structure element mapping rules describe that elements of same + structure variable should allocate/deallocate in a uniform fashion, so + all elements of 's' should be removed together by above 'exit data's. */ + if (d != id) + { + if (omp_target_is_present (&s, d)) + abort (); + if (omp_target_is_present (&s.a, d)) + abort (); + if (omp_target_is_present (&s.b, d)) + abort (); + if (omp_target_is_present (&s.c, d)) + abort (); + if (omp_target_is_present (&s.d, d)) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c new file mode 100644 index 00000000000..9a23b4fbb81 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c @@ -0,0 +1,56 @@ +#include +#include + +struct S +{ + int a, b, c, d, e; +}; +typedef struct S S; + +int main (void) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + S s = { 1, 2, 3, 4, 5 }; + #pragma omp target enter data map (to:s) + + int *p = &s.b; + int *q = &s.d; + #pragma omp target enter data map (alloc: p[:1], q[:1]) + + s.b = 88; + s.d = 99; + + #pragma omp target exit data map (release: s) + if (d != id) + { + if (!omp_target_is_present (&s, d)) + abort (); + if (!omp_target_is_present (&p[0], d)) + abort (); + if (!omp_target_is_present (&q[0], d)) + abort (); + } + + #pragma omp target exit data map (from: q[:1]) + if (d != id) + { + if (omp_target_is_present (&s, d)) + abort (); + if (omp_target_is_present (&p[0], d)) + abort (); + if (omp_target_is_present (&q[0], d)) + abort (); + + if (q[0] != 4) + abort (); + if (p[0] != 88) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c new file mode 100644 index 00000000000..814c30120e5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c @@ -0,0 +1,20 @@ +/* { dg-do run } */ + +struct S +{ + int a, b, c; +}; +typedef struct S S; + +int main (void) +{ + S s; + #pragma omp target data map (alloc: s.a, s.c) + { + #pragma omp target enter data map (alloc: s.b) + } + + return 0; +} +/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) structure element when other mapped elements from the same structure weren't mapped together with it" } */ +/* { dg-shouldfail "" } */