Hi Jakub,
this is a v3 version of my OpenMP 5.0 structure element mapping patch,
v2 was here: https://gcc.gnu.org/pipermail/gcc-patches/2020-December/561139.html
This v3 adds a small bug fix, where the initialization of the refcount didn't
handle all cases, fixed by using gomp_refcount_increment here (more consistent).
I know you had performance concerns in the last round, compared with your
sorting
approach. I'll try to research on that later. Getting the v3 patch posted before
backporting to devel/omp/gcc-11.
Thanks,
Chung-Lin
libgomp/
* 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 --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 405574dfa2b..5cdd110071c 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 cf1baf6291d..ca28924a35b 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..5b7c31406c6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/refcount-1.c
@@ -0,0 +1,52 @@
+#include <omp.h>
+#include <stdlib.h>
+
+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 (!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 (!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 (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..c49d8c12c05
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
@@ -0,0 +1,29 @@
+#include <omp.h>
+#include <stdlib.h>
+
+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 (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..555c6e3e8e0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
@@ -0,0 +1,44 @@
+#include <omp.h>
+#include <stdlib.h>
+
+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 (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..4850eabd879
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
@@ -0,0 +1,63 @@
+#include <omp.h>
+#include <stdlib.h>
+
+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 (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 (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..d50fbf87c02
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c
@@ -0,0 +1,50 @@
+#include <omp.h>
+#include <stdlib.h>
+
+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 (!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 (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 "" } */
diff --git a/libgomp/target.c b/libgomp/target.c
index d373fed8956..f0a66fc83c3 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1559,7 +1559,7 @@ 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 = 1;
+ k->refcount = 0;
k->dynamic_refcount = 0;
if (field_tgt_clear != FIELD_TGT_EMPTY)
{
@@ -1572,7 +1572,6 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
/* Set to first structure element of sequence. */
k->refcount |= REFCOUNT_STRUCTELEM_FLAG_FIRST;
- k->structelem_refcount = 1;
field_tgt_structelem_first = k;
}
else
@@ -1596,6 +1595,11 @@ 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);
