git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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.
This commit is contained in:
Chung-Lin Tang 2021-06-17 21:33:32 +08:00
parent 967b465302
commit 275c736e73
11 changed files with 709 additions and 136 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
libgomp/hashtab.h
View File

@ -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); 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. */ /* Create hash table of size SIZE. */
static htab_t static htab_t
@ -238,11 +247,8 @@ htab_create (size_t size)
result = (htab_t) htab_alloc (sizeof (struct htab) result = (htab_t) htab_alloc (sizeof (struct htab)
+ size * sizeof (hash_entry_type)); + size * sizeof (hash_entry_type));
result->size = size; result->size = size;
result->n_elements = 0;
result->n_deleted = 0;
result->size_prime_index = size_prime_index; result->size_prime_index = size_prime_index;
memset (result->entries, 0, size * sizeof (hash_entry_type)); return htab_clear (result);
return result;
} }
/* Similar to htab_find_slot, but without several unwanted side effects: /* Similar to htab_find_slot, but without several unwanted side effects:

66
libgomp/libgomp.h
View File

@ -1012,11 +1012,35 @@ struct target_mem_desc {
struct target_var_desc list[]; 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. */ /* 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 /* Special value for refcount - tgt_offset contains target address of the
artificial pointer to "omp declare target link" object. */ 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. */ /* Special offset values. */
#define OFFSET_INLINED (~(uintptr_t) 0) #define OFFSET_INLINED (~(uintptr_t) 0)
@ -1044,8 +1068,22 @@ struct splay_tree_key_s {
uintptr_t tgt_offset; uintptr_t tgt_offset;
/* Reference count. */ /* Reference count. */
uintptr_t refcount; uintptr_t refcount;
/* Dynamic reference count. */ union {
uintptr_t dynamic_refcount; /* 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; 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 *, extern void gomp_detach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree_key, struct goacc_asyncqueue *, splay_tree_key,
uintptr_t, bool, struct gomp_coalesce_buf *); uintptr_t, bool, struct gomp_coalesce_buf *);
extern struct target_mem_desc *goacc_map_vars (struct gomp_device_descr *,
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *, struct goacc_asyncqueue *,
size_t, void **, void **, size_t, void **, void **,
size_t *, void *, bool, size_t *, void *, bool,
enum gomp_map_vars_kind); enum gomp_map_vars_kind);
extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *, extern void goacc_unmap_vars (struct target_mem_desc *, bool,
struct goacc_asyncqueue *, 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 void gomp_init_device (struct gomp_device_descr *); extern void gomp_init_device (struct gomp_device_descr *);
extern bool gomp_fini_device (struct gomp_device_descr *); extern bool gomp_fini_device (struct gomp_device_descr *);
extern void gomp_unload_device (struct gomp_device_descr *); extern void gomp_unload_device (struct gomp_device_descr *);

19
libgomp/oacc-mem.c
View File

@ -402,9 +402,8 @@ acc_map_data (void *h, void *d, size_t s)
gomp_mutex_unlock (&acc_dev->lock); gomp_mutex_unlock (&acc_dev->lock);
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes, = goacc_map_vars (acc_dev, NULL, mapnum, &hostaddrs, &devaddrs, &sizes,
&kinds, true, &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
assert (tgt); assert (tgt);
assert (tgt->list_count == 1); assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key; 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); goacc_aq aq = get_goacc_asyncqueue (async);
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
kinds, true, (GOMP_MAP_VARS_OPENACC kinds, true, GOMP_MAP_VARS_ENTER_DATA);
| GOMP_MAP_VARS_ENTER_DATA));
assert (tgt); assert (tgt);
assert (tgt->list_count == 1); assert (tgt->list_count == 1);
n = tgt->list[0].key; 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 /* 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. */ "exit data" called for them at different times. */
static void 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); gomp_mutex_unlock (&acc_dev->lock);
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, = goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true, &sizes[i], &kinds[i], true,
(GOMP_MAP_VARS_OPENACC GOMP_MAP_VARS_ENTER_DATA);
| GOMP_MAP_VARS_ENTER_DATA));
assert (tgt); assert (tgt);
gomp_mutex_lock (&acc_dev->lock); gomp_mutex_lock (&acc_dev->lock);

22
libgomp/oacc-parallel.c
View File

@ -290,8 +290,8 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
goacc_aq aq = get_goacc_asyncqueue (async); goacc_aq aq = get_goacc_asyncqueue (async);
tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, tgt = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
true, GOMP_MAP_VARS_OPENACC); true, 0);
if (profiling_p) if (profiling_p)
{ {
prof_info.event_type = acc_ev_enter_data_end; 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, goacc_profiling_dispatch (&prof_info, &enter_exit_data_event_info,
&api_info); &api_info);
} }
devaddrs = gomp_alloca (sizeof (void *) * mapnum); devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++) for (i = 0; i < mapnum; i++)
devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, 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); &api_info);
} }
/* If running synchronously, unmap immediately. */ /* If running synchronously (aq == NULL), this will unmap immediately. */
if (aq == NULL) goacc_unmap_vars (tgt, true, aq);
gomp_unmap_vars (tgt, true);
else
gomp_unmap_vars_async (tgt, true, aq);
if (profiling_p) if (profiling_p)
{ {
@ -456,8 +453,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
{ {
prof_info.device_type = acc_device_host; prof_info.device_type = acc_device_host;
api_info.device_type = prof_info.device_type; api_info.device_type = prof_info.device_type;
tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, tgt = goacc_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, 0);
GOMP_MAP_VARS_OPENACC);
tgt->prev = thr->mapped_data; tgt->prev = thr->mapped_data;
thr->mapped_data = tgt; 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__); gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, tgt = goacc_map_vars (acc_dev, NULL, mapnum, hostaddrs, NULL, sizes, kinds,
GOMP_MAP_VARS_OPENACC); true, 0);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
tgt->prev = thr->mapped_data; tgt->prev = thr->mapped_data;
thr->mapped_data = tgt; thr->mapped_data = tgt;
@ -542,7 +538,7 @@ GOACC_data_end (void)
gomp_debug (0, " %s: restore mappings\n", __FUNCTION__); gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
thr->mapped_data = tgt->prev; thr->mapped_data = tgt->prev;
gomp_unmap_vars (tgt, true); goacc_unmap_vars (tgt, true, NULL);
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__); gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
if (profiling_p) if (profiling_p)

442
libgomp/target.c
View File

@ -44,6 +44,23 @@
#include "plugin-suffix.h" #include "plugin-suffix.h"
#endif #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) #define FIELD_TGT_EMPTY (~(size_t) 0)
static void gomp_target_init (void); 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 /* Handle the case where gomp_map_lookup, splay_tree_lookup or
gomp_map_0len_lookup found oldn for newn. gomp_map_0len_lookup found oldn for newn.
Helper function of gomp_map_vars. */ 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, struct goacc_asyncqueue *aq, splay_tree_key oldn,
splay_tree_key newn, struct target_var_desc *tgt_var, splay_tree_key newn, struct target_var_desc *tgt_var,
unsigned char kind, bool always_to_flag, 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); assert (kind != GOMP_MAP_ATTACH);
@ -398,8 +523,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
(void *) newn->host_start, (void *) newn->host_start,
newn->host_end - newn->host_start, cbuf); newn->host_end - newn->host_start, cbuf);
if (oldn->refcount != REFCOUNT_INFINITY) gomp_increment_refcount (oldn, refcount_set);
oldn->refcount++;
} }
static int static int
@ -453,7 +577,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
struct goacc_asyncqueue *aq, splay_tree_key n, struct goacc_asyncqueue *aq, splay_tree_key n,
size_t first, size_t i, void **hostaddrs, size_t first, size_t i, void **hostaddrs,
size_t *sizes, void *kinds, 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 gomp_device_descr *devicep = tgt->device_descr;
struct splay_tree_s *mem_map = &devicep->mem_map; 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) && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{ {
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
kind & typemask, false, cbuf); kind & typemask, false, cbuf, refcount_set);
return; return;
} }
if (sizes[i] == 0) if (sizes[i] == 0)
@ -487,7 +611,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
== n2->tgt_offset - n->tgt_offset) == n2->tgt_offset - n->tgt_offset)
{ {
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
kind & typemask, false, cbuf); kind & typemask, false, cbuf, refcount_set);
return; 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) && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{ {
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
kind & typemask, false, cbuf); kind & typemask, false, cbuf, refcount_set);
return; return;
} }
} }
@ -671,11 +795,13 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum, struct goacc_asyncqueue *aq, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void **hostaddrs, void **devaddrs, size_t *sizes,
void *kinds, bool short_mapkind, void *kinds, bool short_mapkind,
htab_t *refcount_set,
enum gomp_map_vars_kind pragma_kind) enum gomp_map_vars_kind pragma_kind)
{ {
size_t i, tgt_align, tgt_size, not_found_cnt = 0; size_t i, tgt_align, tgt_size, not_found_cnt = 0;
bool has_firstprivate = false; bool has_firstprivate = false;
bool has_always_ptrset = 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 rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7; const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map; 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++) for (i = first; i <= last; i++)
gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs, gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs,
sizes, kinds, NULL); sizes, kinds, NULL, refcount_set);
i--; i--;
continue; 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], 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; i += always_to_cnt;
} }
else else
@ -1022,6 +1149,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_node array = tgt->array; splay_tree_node array = tgt->array;
size_t j, field_tgt_offset = 0, field_tgt_clear = FIELD_TGT_EMPTY; size_t j, field_tgt_offset = 0, field_tgt_clear = FIELD_TGT_EMPTY;
uintptr_t field_tgt_base = 0; uintptr_t field_tgt_base = 0;
splay_tree_key field_tgt_structelem_first = NULL;
for (i = 0; i < mapnum; i++) for (i = 0; i < mapnum; i++)
if (has_always_ptrset 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].copy_from = false;
tgt->list[j].always_copy_from = false; tgt->list[j].always_copy_from = false;
tgt->list[j].is_attach = false; tgt->list[j].is_attach = false;
if (k->refcount != REFCOUNT_INFINITY) gomp_increment_refcount (k, refcount_set);
k->refcount++;
gomp_map_pointer (k->tgt, aq, gomp_map_pointer (k->tgt, aq,
(uintptr_t) *(void **) hostaddrs[j], (uintptr_t) *(void **) hostaddrs[j],
k->tgt_offset + ((uintptr_t) 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_base = (uintptr_t) hostaddrs[first];
field_tgt_offset = tgt_size; field_tgt_offset = tgt_size;
field_tgt_clear = last; field_tgt_clear = last;
field_tgt_structelem_first = NULL;
tgt_size += cur_node.host_end tgt_size += cur_node.host_end
- (uintptr_t) hostaddrs[first]; - (uintptr_t) hostaddrs[first];
continue; continue;
} }
for (i = first; i <= last; i++) for (i = first; i <= last; i++)
gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs, gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs,
sizes, kinds, cbufp); sizes, kinds, cbufp, refcount_set);
i--; i--;
continue; continue;
case GOMP_MAP_ALWAYS_POINTER: 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); splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n && n->refcount != REFCOUNT_LINK) if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
kind & typemask, false, cbufp); kind & typemask, false, cbufp,
refcount_set);
else else
{ {
k->aux = NULL; k->aux = NULL;
@ -1252,10 +1381,33 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
size_t align = (size_t) 1 << (kind >> rshift); size_t align = (size_t) 1 << (kind >> rshift);
tgt->list[i].key = k; tgt->list[i].key = k;
k->tgt = tgt; k->tgt = tgt;
k->refcount = 0;
k->dynamic_refcount = 0;
if (field_tgt_clear != FIELD_TGT_EMPTY) if (field_tgt_clear != FIELD_TGT_EMPTY)
{ {
k->tgt_offset = k->host_start - field_tgt_base k->tgt_offset = k->host_start - field_tgt_base
+ field_tgt_offset; + 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) if (i == field_tgt_clear)
field_tgt_clear = FIELD_TGT_EMPTY; field_tgt_clear = FIELD_TGT_EMPTY;
} }
@ -1265,14 +1417,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
k->tgt_offset = tgt_size; k->tgt_offset = tgt_size;
tgt_size += k->host_end - k->host_start; 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].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
tgt->list[i].always_copy_from tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask); = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
tgt->list[i].is_attach = false; tgt->list[i].is_attach = false;
tgt->list[i].offset = 0; tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start; tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
k->dynamic_refcount = 0;
tgt->refcount++; tgt->refcount++;
array->left = NULL; array->left = NULL;
array->right = 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].always_copy_from = false;
tgt->list[j].is_attach = false; tgt->list[j].is_attach = false;
tgt->list[i].has_null_ptr_assoc |= !(*(void **) hostaddrs[j]); tgt->list[i].has_null_ptr_assoc |= !(*(void **) hostaddrs[j]);
if (k->refcount != REFCOUNT_INFINITY) /* For OpenMP, the use of refcount_sets causes
k->refcount++; 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, gomp_map_pointer (tgt, aq,
(uintptr_t) *(void **) hostaddrs[j], (uintptr_t) *(void **) hostaddrs[j],
k->tgt_offset k->tgt_offset
@ -1426,24 +1587,41 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
return tgt; return tgt;
} }
attribute_hidden struct target_mem_desc * static struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, 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, /* This management of a local refcount_set is for convenience of callers
sizes, kinds, short_mapkind, pragma_kind); 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 * attribute_hidden struct target_mem_desc *
gomp_map_vars_async (struct gomp_device_descr *devicep, goacc_map_vars (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum, struct goacc_asyncqueue *aq, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void **hostaddrs, void **devaddrs, size_t *sizes,
void *kinds, bool short_mapkind, void *kinds, bool short_mapkind,
enum gomp_map_vars_kind pragma_kind) enum gomp_map_vars_kind pragma_kind)
{ {
return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs, 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 static void
@ -1481,22 +1659,56 @@ gomp_unref_tgt_void (void *ptr)
(void) gomp_unref_tgt (ptr); (void) gomp_unref_tgt (ptr);
} }
static inline __attribute__((always_inline)) bool static void
gomp_remove_var_internal (struct gomp_device_descr *devicep, splay_tree_key k, gomp_remove_splay_tree_key (splay_tree sp, splay_tree_key k)
struct goacc_asyncqueue *aq)
{ {
bool is_tgt_unmapped = false; splay_tree_remove (sp, k);
splay_tree_remove (&devicep->mem_map, k);
if (k->aux) if (k->aux)
{ {
if (k->aux->link_key) if (k->aux->link_key)
splay_tree_insert (&devicep->mem_map, splay_tree_insert (sp, (splay_tree_node) k->aux->link_key);
(splay_tree_node) k->aux->link_key);
if (k->aux->attach_count) if (k->aux->attach_count)
free (k->aux->attach_count); free (k->aux->attach_count);
free (k->aux); free (k->aux);
k->aux = NULL; 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) if (aq)
devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void, devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
(void *) k->tgt); (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 static inline __attribute__((always_inline)) void
gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, 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; 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) if (tgt->list[i].is_attach)
continue; continue;
bool do_unmap = false; bool do_copy, do_remove;
if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) gomp_decrement_refcount (k, refcount_set, false, &do_copy, &do_remove);
k->refcount--;
else if (k->refcount == 1)
{
k->refcount--;
do_unmap = true;
}
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) || tgt->list[i].always_copy_from)
gomp_copy_dev2host (devicep, aq, gomp_copy_dev2host (devicep, aq,
(void *) (k->host_start + tgt->list[i].offset), (void *) (k->host_start + tgt->list[i].offset),
(void *) (k->tgt->tgt_start + k->tgt_offset (void *) (k->tgt->tgt_start + k->tgt_offset
+ tgt->list[i].offset), + tgt->list[i].offset),
tgt->list[i].length); tgt->list[i].length);
if (do_unmap) if (do_remove)
{ {
struct target_mem_desc *k_tgt = k->tgt; struct target_mem_desc *k_tgt = k->tgt;
bool is_tgt_unmapped = gomp_remove_var (devicep, k); 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); gomp_mutex_unlock (&devicep->lock);
} }
attribute_hidden void static void
gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom) 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 attribute_hidden void
gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom, goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom,
struct goacc_asyncqueue *aq) struct goacc_asyncqueue *aq)
{ {
gomp_unmap_vars_internal (tgt, do_copyfrom, aq); gomp_unmap_vars_internal (tgt, do_copyfrom, NULL, aq);
} }
static void 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))) || !(fn_addr = gomp_get_target_fn_addr (devicep, fn)))
return gomp_target_fallback (fn, hostaddrs, devicep); return gomp_target_fallback (fn, hostaddrs, devicep);
htab_t refcount_set = htab_create (mapnum);
struct target_mem_desc *tgt_vars struct target_mem_desc *tgt_vars
= gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, = 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, devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start,
NULL); 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 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; struct target_mem_desc *tgt_vars;
htab_t refcount_set = NULL;
if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
{ {
if (!fpc_done) if (!fpc_done)
@ -2285,13 +2509,21 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
tgt_vars = NULL; tgt_vars = NULL;
} }
else 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, devicep->run_func (devicep->target_id, fn_addr,
tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs, tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
args); args);
if (tgt_vars) 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. */ /* 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. */ would get out of sync. */
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false,
GOMP_MAP_VARS_DATA); NULL, GOMP_MAP_VARS_DATA);
tgt->prev = icv->target_data; tgt->prev = icv->target_data;
icv->target_data = tgt; icv->target_data = tgt;
} }
@ -2333,7 +2565,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false, = 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); struct gomp_task_icv *icv = gomp_icv (true);
tgt->prev = icv->target_data; tgt->prev = icv->target_data;
icv->target_data = tgt; icv->target_data = tgt;
@ -2352,7 +2584,7 @@ GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
struct target_mem_desc *tgt struct target_mem_desc *tgt
= gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true, = 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); struct gomp_task_icv *icv = gomp_icv (true);
tgt->prev = icv->target_data; tgt->prev = icv->target_data;
icv->target_data = tgt; icv->target_data = tgt;
@ -2366,7 +2598,7 @@ GOMP_target_end_data (void)
{ {
struct target_mem_desc *tgt = icv->target_data; struct target_mem_desc *tgt = icv->target_data;
icv->target_data = tgt->prev; 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 static void
gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, 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; const int typemask = 0xff;
size_t i; size_t i;
@ -2489,6 +2722,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
false, NULL); false, NULL);
} }
int nrmvars = 0;
splay_tree_key remove_vars[mapnum];
for (i = 0; i < mapnum; i++) for (i = 0; i < mapnum; i++)
{ {
struct splay_tree_key_s cur_node; struct splay_tree_key_s cur_node;
@ -2510,22 +2746,32 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
if (!k) if (!k)
continue; continue;
if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY) bool delete_p = (kind == GOMP_MAP_DELETE
k->refcount--; || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION);
if ((kind == GOMP_MAP_DELETE bool do_copy, do_remove;
|| kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION) gomp_decrement_refcount (k, refcount_set, delete_p, &do_copy,
&& k->refcount != REFCOUNT_INFINITY) &do_remove);
k->refcount = 0;
if ((kind == GOMP_MAP_FROM && k->refcount == 0) if ((kind == GOMP_MAP_FROM && do_copy)
|| kind == GOMP_MAP_ALWAYS_FROM) || kind == GOMP_MAP_ALWAYS_FROM)
gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start, gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
(void *) (k->tgt->tgt_start + k->tgt_offset (void *) (k->tgt->tgt_start + k->tgt_offset
+ cur_node.host_start + cur_node.host_start
- k->host_start), - k->host_start),
cur_node.host_end - cur_node.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; break;
case GOMP_MAP_DETACH: 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); 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 /* The variables are mapped separately such that they can be released
independently. */ independently. */
size_t i, j; 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) if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
{ {
gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i], 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]; i += sizes[i];
} }
else if ((kinds[i] & 0xff) == GOMP_MAP_TO_PSET) 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)) && !GOMP_MAP_ALWAYS_POINTER_P (get_kind (true, kinds, j) & 0xff))
break; break;
gomp_map_vars (devicep, j-i, &hostaddrs[i], NULL, &sizes[i], 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; i += j - i - 1;
} }
else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH) 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 /* An attach operation must be processed together with the mapped
base-pointer list item. */ base-pointer list item. */
gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i], 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; i += 1;
} }
else else
gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i], 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 else
gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds); gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds, &refcount_set);
htab_free (refcount_set);
} }
bool bool
@ -2674,7 +2928,7 @@ gomp_target_task_fn (void *data)
if (ttask->state == GOMP_TARGET_TASK_FINISHED) if (ttask->state == GOMP_TARGET_TASK_FINISHED)
{ {
if (ttask->tgt) if (ttask->tgt)
gomp_unmap_vars (ttask->tgt, true); gomp_unmap_vars (ttask->tgt, true, NULL);
return false; return false;
} }
@ -2688,7 +2942,7 @@ gomp_target_task_fn (void *data)
{ {
ttask->tgt = gomp_map_vars (devicep, ttask->mapnum, ttask->hostaddrs, ttask->tgt = gomp_map_vars (devicep, ttask->mapnum, ttask->hostaddrs,
NULL, ttask->sizes, ttask->kinds, true, NULL, ttask->sizes, ttask->kinds, true,
GOMP_MAP_VARS_TARGET); NULL, GOMP_MAP_VARS_TARGET);
actual_arguments = (void *) ttask->tgt->tgt_start; actual_arguments = (void *) ttask->tgt->tgt_start;
} }
ttask->state = GOMP_TARGET_TASK_READY_TO_RUN; 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) if (ttask->flags & GOMP_TARGET_FLAG_UPDATE)
gomp_update (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes, gomp_update (devicep, ttask->mapnum, ttask->hostaddrs, ttask->sizes,
ttask->kinds, true); 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 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; return false;
} }

61
libgomp/testsuite/libgomp.c-c++-common/refcount-1.c Normal file
View File

@ -0,0 +1,61 @@
#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 (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;
}

29
libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c Normal file
View File

@ -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 (d != id && omp_target_is_present (&s.a, d))
abort ();
return 0;
}

47
libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c Normal file
View File

@ -0,0 +1,47 @@
#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 (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;
}

69
libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c Normal file
View File

@ -0,0 +1,69 @@
#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 (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;
}

56
libgomp/testsuite/libgomp.c-c++-common/struct-elem-4.c Normal file
View File

@ -0,0 +1,56 @@
#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 (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;
}

20
libgomp/testsuite/libgomp.c-c++-common/struct-elem-5.c Normal file
View File

@ -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 "" } */