On Sat, 9 Nov 2019 01:28:51 +0000
Julian Brown <jul...@codesourcery.com> wrote:
> On Thu, 31 Oct 2019 19:11:57 +0100
> Thomas Schwinge <tho...@codesourcery.com> wrote:
>
> > So that's not related to reference counting, needs to be discussed
> > separately.
> >
> > ..., and while I do agree that the current code is a bit "strange"
> > (returning 'tgt->to_free'), I couldn't quickly find or come up with
> > a test cases where this would actually do the wrong thing. After
> > all, this is the code path taken for "not present", and 'tgt' is
> > built anew for one single mapping, with no alignment set (which
> > would cause 'to_free' to differ from 'tgt_start'); 'tgt_offset'
> > should always be zero, and 'h' always the same as 'host_start'.
> > What am I missing? That is, given the current set of libgomp test
> > cases, the attached never triggeres.
>
> The code can't stay exactly as it is with this patch, because the tgt
> return value from gomp_map_vars_async with
> GOMP_MAP_VARS_OPENACC_ENTER_DATA is a null pointer.
>
> So, the device pointer calculation needed to be re-done -- although
> it's not quite a bug fix, as you point out, and some of the offsets
> will always be zero or cancel out in practice.
>
> *However*, it looks like the device pointer calculation for the
> "present" case is wrong in the preceding code. I've addressed that in
> the patch posted here:
>
> https://gcc.gnu.org/ml/gcc-patches/2019-11/msg00661.html
>
> The patch attached here applies on top of that one, and attempts to
> keep the device pointer calculation "the same" for the non-present
> case, modulo an extra lookup_host -- and also adds some assertions to
> make sure the assumptions about zero/cancelled-out offsets stay true.
Here's another iteration that applies over the version of the
present/subarray patch committed, and also addresses the use of
REFCOUNT_INFINITY on target blocks as queried in the following message:
https://gcc.gnu.org/ml/gcc-patches/2019-11/msg01146.html
Most uses of REFCOUNT_INFINITY indeed appear to be unreachable (as in,
a target_mem_desc with refcount == REFCOUNT_INFINITY will most of the
time be linked from a splay tree key with refcount ==
REFCOUNT_INFINITY, and the code to decrement the former's refcount
and/or free the block will never be called).
I found one case (for OpenACC) where a runtime check/error can be
added -- attempting to free a mapped target block corresponding to a
device_resident global variable using an API routine. I don't think
there's a code path using directives (for either OpenACC or OpenMP) that
exhibits any problematic behaviour in that regard. I've added a couple
of test cases, and a couple of assertions.
OK now? (Or perhaps the REFCOUNT_INFINITY bits want splitting out? It
all still arguably comes under the "refcount overhaul" umbrella!).
Thanks,
Julian
commit d027f9ab63fdf871076f40e2cf42f672ed3f0e69
Author: Julian Brown <jul...@codesourcery.com>
Date: Mon Nov 5 15:51:46 2018 -0800
OpenACC reference count overhaul
2019-11-22 Julian Brown <jul...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
libgomp/
* libgomp.h (struct splay_tree_key_s): Substitute dynamic_refcount
field for virtual_refcount.
(struct acc_dispatch_t): Remove data_environ field.
(struct gomp_device_descr): Update comment on openacc field.
(enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA.
(gomp_acc_insert_pointer, gomp_acc_remove_pointer, gomp_free_memmap):
Remove prototypes.
(gomp_remove_var_async): Add prototype.
* oacc-host.c (host_dispatch): Don't initialise removed data_environ
field.
* oacc-init.c (acc_shutdown_1): Iteratively call gomp_remove_var
instead of calling gomp_free_memmap.
* oacc-mem.c (lookup_dev_1): New function.
(lookup_dev): Reimplement using above.
(acc_free, acc_hostptr): Update calls to lookup_dev.
(acc_map_data): Likewise. Don't add to data_environ list.
(acc_unmap_data): Remove call to gomp_unmap_vars. Fix semantics to
remove mapping, but not mapped data. Handle REFCOUNT_INFINITY on
target blocks.
(present_create_copy): Use virtual_refcount instead of
dynamic_refcount. Don't manipulate data_environ. Re-do lookup for
target pointer return value.
(delete_copyout): Update for virtual_refcount semantics. Use
goacc_remove_var_async for asynchronous delete/copyouts.
(gomp_acc_insert_pointer, gomp_acc_remove_pointer): Remove functions.
* oacc-parallel.c (find_pointer): Remove function.
(find_group_last, goacc_enter_data_internal,
goacc_exit_data_internal): New functions.
(GOACC_enter_exit_data): Use goacc_enter_data_internal and
goacc_exit_data_internal helper functions.
* target.c (gomp_map_vars_internal): Handle
GOMP_MAP_VARS_OPENACC_ENTER_DATA. Update for virtual_refcount
semantics.
(gomp_unref_tgt): Add assertion for tgt->refcount not being
REFCOUNT_INFINITY.
(gomp_unref_tgt_void): New function.
(gomp_remove_var): Reimplement in terms of...
(gomp_remove_var_internal): ...this new helper function.
(gomp_remove_var_async): Implement using above helper function.
(gomp_unmap_vars_internal): Update for virtual_refcount
semantics.
(gomp_load_image_to_device): Zero-initialise virtual_refcount field.
(gomp_free_memmap): Remove function.
(omp_target_associate_ptr): Zero-initialise virtual_refcount and
link_key splay tree key fields.
(gomp_exit_data): Add assertion for tgt->refcount not being
REFCOUNT_INFINITY.
(gomp_target_init): Don't initialise removed data_environ field.
* testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c: New test.
* testsuite/libgomp.c-c++-common/unmap-infinity-2.c: New test.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index bab733d2b2d..521e0803122 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1000,8 +1000,11 @@ struct splay_tree_key_s {
uintptr_t tgt_offset;
/* Reference count. */
uintptr_t refcount;
- /* Dynamic reference count. */
- uintptr_t dynamic_refcount;
+ /* Reference counts beyond those that represent genuine references in the
+ linked splay tree key/target memory structures, e.g. for multiple OpenACC
+ "present increment" operations (via "acc enter data") referring to the same
+ host-memory block. */
+ uintptr_t virtual_refcount;
/* Pointer to the original mapping of "omp declare target link" object. */
splay_tree_key link_key;
};
@@ -1025,13 +1028,6 @@ splay_compare (splay_tree_key x, splay_tree_key y)
typedef struct acc_dispatch_t
{
- /* This is a linked list of data mapped using the
- acc_map_data/acc_unmap_data or "acc enter data"/"acc exit data" pragmas.
- Unlike mapped_data in the goacc_thread struct, unmapping can
- happen out-of-order with respect to mapping. */
- /* This is guarded by the lock in the "outer" struct gomp_device_descr. */
- struct target_mem_desc *data_environ;
-
/* Execute. */
__typeof (GOMP_OFFLOAD_openacc_exec) *exec_func;
@@ -1132,8 +1128,7 @@ struct gomp_device_descr
enum gomp_device_state state;
/* OpenACC-specific data and functions. */
- /* This is mutable because of its mutable data_environ and target_data
- members. */
+ /* This is mutable because of the target_data member. */
acc_dispatch_t openacc;
};
@@ -1141,13 +1136,12 @@ struct gomp_device_descr
enum gomp_map_vars_kind
{
GOMP_MAP_VARS_OPENACC,
+ GOMP_MAP_VARS_OPENACC_ENTER_DATA,
GOMP_MAP_VARS_TARGET,
GOMP_MAP_VARS_DATA,
GOMP_MAP_VARS_ENTER_DATA
};
-extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *, int);
-extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
unsigned short *);
struct gomp_coalesce_buf;
@@ -1173,9 +1167,10 @@ extern void gomp_unmap_vars_async (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_free_memmap (struct splay_tree_s *);
extern void gomp_unload_device (struct gomp_device_descr *);
extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key);
+extern void gomp_remove_var_async (struct gomp_device_descr *, splay_tree_key,
+ struct goacc_asyncqueue *);
/* work.c */
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index cbcac9bf7b3..e9cd4bfcd4a 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -264,8 +264,6 @@ static struct gomp_device_descr host_dispatch =
.state = GOMP_DEVICE_UNINITIALIZED,
.openacc = {
- .data_environ = NULL,
-
.exec_func = host_openacc_exec,
.create_thread_data_func = host_openacc_create_thread_data,
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index e1568c535b3..8465ed503a3 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -357,7 +357,14 @@ acc_shutdown_1 (acc_device_t d)
if (walk->dev)
{
gomp_mutex_lock (&walk->dev->lock);
- gomp_free_memmap (&walk->dev->mem_map);
+
+ while (walk->dev->mem_map.root)
+ {
+ splay_tree_key k = &walk->dev->mem_map.root->key;
+ k->link_key = NULL;
+ gomp_remove_var (walk->dev, k);
+ }
+
gomp_mutex_unlock (&walk->dev->lock);
walk->dev = NULL;
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index aafe88d3a14..34795437b21 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -50,42 +50,39 @@ lookup_host (struct gomp_device_descr *dev, void *h, size_t s)
return key;
}
-/* Return block containing [D->S), or NULL if not contained.
- The list isn't ordered by device address, so we have to iterate
- over the whole array. This is not expected to be a common
- operation. The device lock associated with TGT must be locked on entry, and
- remains locked on exit. */
+/* Helper for lookup_dev. Iterate over splay tree. */
static splay_tree_key
-lookup_dev (struct target_mem_desc *tgt, void *d, size_t s)
+lookup_dev_1 (splay_tree_node node, uintptr_t d, size_t s)
{
- int i;
- struct target_mem_desc *t;
+ splay_tree_key k = &node->key, found = NULL;
+ struct target_mem_desc *t = k->tgt;
- if (!tgt)
- return NULL;
+ if (d >= t->tgt_start && d + s <= t->tgt_end)
+ return k;
- for (t = tgt; t != NULL; t = t->prev)
- {
- if (t->tgt_start <= (uintptr_t) d && t->tgt_end >= (uintptr_t) d + s)
- break;
- }
+ if (node->left)
+ found = lookup_dev_1 (node->left, d, s);
- if (!t)
- return NULL;
+ if (!found && node->right)
+ found = lookup_dev_1 (node->right, d, s);
- for (i = 0; i < t->list_count; i++)
- {
- void * offset;
+ return found;
+}
- splay_tree_key k = &t->array[i].key;
- offset = d - t->tgt_start + k->tgt_offset;
+/* Return block containing [D->S), or NULL if not contained.
+ The list isn't ordered by device address, so we have to iterate
+ over the whole array. This is not expected to be a common
+ operation. The device lock associated with TGT must be locked on entry, and
+ remains locked on exit. */
- if (k->host_start + offset <= (void *) k->host_end)
- return k;
- }
+static splay_tree_key
+lookup_dev (splay_tree mem_map, void *d, size_t s)
+{
+ if (!mem_map || !mem_map->root)
+ return NULL;
- return NULL;
+ return lookup_dev_1 (mem_map->root, (uintptr_t) d, s);
}
/* OpenACC is silent on how memory exhaustion is indicated. We return
@@ -150,7 +147,7 @@ acc_free (void *d)
/* We don't have to call lazy open here, as the ptr value must have
been returned by acc_malloc. It's not permitted to pass NULL in
(unless you got that null from acc_malloc). */
- if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1)))
+ if ((k = lookup_dev (&acc_dev->mem_map, d, 1)))
{
void *offset;
@@ -301,7 +298,7 @@ acc_hostptr (void *d)
gomp_mutex_lock (&acc_dev->lock);
- n = lookup_dev (acc_dev->openacc.data_environ, d, 1);
+ n = lookup_dev (&acc_dev->mem_map, d, 1);
if (!n)
{
@@ -396,7 +393,7 @@ acc_map_data (void *h, void *d, size_t s)
(int)s);
}
- if (lookup_dev (thr->dev->openacc.data_environ, d, s))
+ if (lookup_dev (&thr->dev->mem_map, d, s))
{
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("device address [%p, +%d] is already mapped", (void *)d,
@@ -415,11 +412,6 @@ acc_map_data (void *h, void *d, size_t s)
thr->api_info = NULL;
}
}
-
- gomp_mutex_lock (&acc_dev->lock);
- tgt->prev = acc_dev->openacc.data_environ;
- acc_dev->openacc.data_environ = tgt;
- gomp_mutex_unlock (&acc_dev->lock);
}
void
@@ -438,12 +430,9 @@ acc_unmap_data (void *h)
acc_api_info api_info;
bool profiling_p = GOACC_PROFILING_SETUP_P (thr, &prof_info, &api_info);
- size_t host_size;
-
gomp_mutex_lock (&acc_dev->lock);
splay_tree_key n = lookup_host (acc_dev, h, 1);
- struct target_mem_desc *t;
if (!n)
{
@@ -451,47 +440,33 @@ acc_unmap_data (void *h)
gomp_fatal ("%p is not a mapped block", (void *)h);
}
- host_size = n->host_end - n->host_start;
-
if (n->host_start != (uintptr_t) h)
{
+ size_t host_size = n->host_end - n->host_start;
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("[%p,%d] surrounds %p",
(void *) n->host_start, (int) host_size, (void *) h);
}
- /* Mark for removal. */
- n->refcount = 1;
+ splay_tree_remove (&acc_dev->mem_map, n);
- t = n->tgt;
+ struct target_mem_desc *tgt = n->tgt;
- if (t->refcount == 2)
+ if (tgt->refcount == REFCOUNT_INFINITY)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("cannot unmap target block");
+ }
+ else if (tgt->refcount > 1)
+ tgt->refcount--;
+ else
{
- struct target_mem_desc *tp;
-
- /* This is the last reference, so pull the descriptor off the
- chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from
- freeing the device memory. */
- t->tgt_end = 0;
- t->to_free = 0;
-
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
-
- break;
- }
+ free (tgt->array);
+ free (tgt);
}
gomp_mutex_unlock (&acc_dev->lock);
- gomp_unmap_vars (t, true);
-
if (profiling_p)
{
thr->prof_info = NULL;
@@ -552,8 +527,9 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
if (n->refcount != REFCOUNT_INFINITY)
{
n->refcount++;
- n->dynamic_refcount++;
+ n->virtual_refcount++;
}
+
gomp_mutex_unlock (&acc_dev->lock);
}
else if (!(f & FLAG_CREATE))
@@ -563,7 +539,6 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
}
else
{
- struct target_mem_desc *tgt;
size_t mapnum = 1;
unsigned short kinds;
void *hostaddrs = h;
@@ -577,17 +552,15 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
goacc_aq aq = get_goacc_asyncqueue (async);
- tgt = gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s,
- &kinds, true, GOMP_MAP_VARS_OPENACC);
- /* Initialize dynamic refcount. */
- tgt->list[0].key->dynamic_refcount = 1;
+ gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s, &kinds,
+ true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
gomp_mutex_lock (&acc_dev->lock);
-
- d = tgt->to_free;
- tgt->prev = acc_dev->openacc.data_environ;
- acc_dev->openacc.data_environ = tgt;
-
+ n = lookup_host (acc_dev, h, s);
+ assert (n != NULL);
+ assert (n->tgt_offset == 0);
+ assert ((uintptr_t) h == n->host_start);
+ d = (void *) n->tgt->tgt_start;
gomp_mutex_unlock (&acc_dev->lock);
}
@@ -670,7 +643,6 @@ static void
delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
{
splay_tree_key n;
- void *d;
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
@@ -699,9 +671,6 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
}
- d = (void *) (n->tgt->tgt_start + n->tgt_offset
- + (uintptr_t) h - n->host_start);
-
if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end)
{
size_t host_size = n->host_end - n->host_start;
@@ -713,48 +682,34 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
if (n->refcount == REFCOUNT_INFINITY)
{
n->refcount = 0;
- n->dynamic_refcount = 0;
- }
- if (n->refcount < n->dynamic_refcount)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
+ n->virtual_refcount = 0;
}
if (f & FLAG_FINALIZE)
{
- n->refcount -= n->dynamic_refcount;
- n->dynamic_refcount = 0;
+ n->refcount -= n->virtual_refcount;
+ n->virtual_refcount = 0;
}
- else if (n->dynamic_refcount)
+
+ if (n->virtual_refcount > 0)
{
- n->dynamic_refcount--;
n->refcount--;
+ n->virtual_refcount--;
}
+ else if (n->refcount > 0)
+ n->refcount--;
if (n->refcount == 0)
{
- if (n->tgt->refcount == 2)
- {
- struct target_mem_desc *tp, *t;
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
- break;
- }
- }
+ goacc_aq aq = get_goacc_asyncqueue (async);
if (f & FLAG_COPYOUT)
{
- goacc_aq aq = get_goacc_asyncqueue (async);
+ void *d = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + (uintptr_t) h - n->host_start);
gomp_copy_dev2host (acc_dev, aq, h, d, s);
}
- gomp_remove_var (acc_dev, n);
+ gomp_remove_var_async (acc_dev, n, aq);
}
gomp_mutex_unlock (&acc_dev->lock);
@@ -890,142 +845,3 @@ acc_update_self_async (void *h, size_t s, int async)
{
update_dev_host (0, h, s, async);
}
-
-void
-gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
- void *kinds, int async)
-{
- struct target_mem_desc *tgt;
- struct goacc_thread *thr = goacc_thread ();
- struct gomp_device_descr *acc_dev = thr->dev;
-
- if (acc_is_present (*hostaddrs, *sizes))
- {
- splay_tree_key n;
- gomp_mutex_lock (&acc_dev->lock);
- n = lookup_host (acc_dev, *hostaddrs, *sizes);
- gomp_mutex_unlock (&acc_dev->lock);
-
- tgt = n->tgt;
- for (size_t i = 0; i < tgt->list_count; i++)
- if (tgt->list[i].key == n)
- {
- for (size_t j = 0; j < mapnum; j++)
- if (i + j < tgt->list_count && tgt->list[i + j].key)
- {
- tgt->list[i + j].key->refcount++;
- tgt->list[i + j].key->dynamic_refcount++;
- }
- return;
- }
- /* Should not reach here. */
- gomp_fatal ("Dynamic refcount incrementing failed for pointer/pset");
- }
-
- gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
- 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);
- gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
-
- /* Initialize dynamic refcount. */
- tgt->list[0].key->dynamic_refcount = 1;
-
- gomp_mutex_lock (&acc_dev->lock);
- tgt->prev = acc_dev->openacc.data_environ;
- acc_dev->openacc.data_environ = tgt;
- gomp_mutex_unlock (&acc_dev->lock);
-}
-
-void
-gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
- int finalize, int mapnum)
-{
- struct goacc_thread *thr = goacc_thread ();
- struct gomp_device_descr *acc_dev = thr->dev;
- splay_tree_key n;
- struct target_mem_desc *t;
- int minrefs = (mapnum == 1) ? 2 : 3;
-
- if (!acc_is_present (h, s))
- return;
-
- gomp_mutex_lock (&acc_dev->lock);
-
- n = lookup_host (acc_dev, h, 1);
-
- if (!n)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("%p is not a mapped block", (void *)h);
- }
-
- gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
-
- t = n->tgt;
-
- if (n->refcount < n->dynamic_refcount)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
- }
-
- if (finalize)
- {
- n->refcount -= n->dynamic_refcount;
- n->dynamic_refcount = 0;
- }
- else if (n->dynamic_refcount)
- {
- n->dynamic_refcount--;
- n->refcount--;
- }
-
- gomp_mutex_unlock (&acc_dev->lock);
-
- if (n->refcount == 0)
- {
- if (t->refcount == minrefs)
- {
- /* This is the last reference, so pull the descriptor off the
- chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from
- freeing the device memory. */
- struct target_mem_desc *tp;
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
- {
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
- break;
- }
- }
- }
-
- /* Set refcount to 1 to allow gomp_unmap_vars to unmap it. */
- n->refcount = 1;
- t->refcount = minrefs;
- for (size_t i = 0; i < t->list_count; i++)
- if (t->list[i].key == n)
- {
- t->list[i].copy_from = force_copyfrom ? 1 : 0;
- break;
- }
-
- /* If running synchronously, unmap immediately. */
- if (async < acc_async_noval)
- gomp_unmap_vars (t, true);
- else
- {
- goacc_aq aq = get_goacc_asyncqueue (async);
- gomp_unmap_vars_async (t, true, aq);
- }
- }
-
- gomp_mutex_unlock (&acc_dev->lock);
-
- gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
-}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 68a60de24fa..87f9ae01906 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -47,23 +47,39 @@ _Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_HOST_FALLBACK)
"legacy GOMP_DEVICE_HOST_FALLBACK broken");
-/* Returns the number of mappings associated with the pointer or pset. PSET
- have three mappings, whereas pointer have two. */
+/* Some types of (pointer) variables use several consecutive mappings, which
+ must be treated as a group for enter/exit data directives. This function
+ returns the last mapping in such a group (inclusive), or POS for singleton
+ mappings. */
static int
-find_pointer (int pos, size_t mapnum, unsigned short *kinds)
+find_group_last (int pos, size_t mapnum, unsigned short *kinds)
{
- if (pos + 1 >= mapnum)
- return 0;
+ unsigned char kind0 = kinds[pos] & 0xff;
+ int first_pos = pos, last_pos = pos;
- unsigned char kind = kinds[pos+1] & 0xff;
-
- if (kind == GOMP_MAP_TO_PSET)
- return 3;
- else if (kind == GOMP_MAP_POINTER)
- return 2;
+ if (kind0 == GOMP_MAP_TO_PSET)
+ {
+ while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
+ last_pos = ++pos;
+ /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */
+ assert (last_pos > first_pos);
+ }
+ else
+ {
+ /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other
+ mapping. */
+ if (pos + 1 < mapnum
+ && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)
+ return pos + 1;
+
+ /* We can have one or several GOMP_MAP_POINTER mappings after a to/from
+ (etc.) mapping. */
+ while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
+ last_pos = ++pos;
+ }
- return 0;
+ return last_pos;
}
/* Handle the mapping pair that are presented when a
@@ -580,6 +596,108 @@ GOACC_data_end (void)
}
}
+/* Map variables for OpenACC "enter data". We can't just call
+ gomp_map_vars_async once, because individual mapped variables might have
+ "exit data" called for them at different times. */
+
+static void
+goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
+ void **hostaddrs, size_t *sizes,
+ unsigned short *kinds, goacc_aq aq)
+{
+ for (size_t i = 0; i < mapnum; i++)
+ {
+ int group_last = find_group_last (i, mapnum, kinds);
+
+ gomp_map_vars_async (acc_dev, aq,
+ (group_last - i) + 1,
+ &hostaddrs[i], NULL,
+ &sizes[i], &kinds[i], true,
+ GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+
+ i = group_last;
+ }
+}
+
+/* Unmap variables for OpenACC "exit data", with optional finalization
+ (affecting all mappings in this operation). */
+
+static void
+goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
+ void **hostaddrs, size_t *sizes,
+ unsigned short *kinds, bool finalize, goacc_aq aq)
+{
+ gomp_mutex_lock (&acc_dev->lock);
+
+ for (size_t i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+ bool copyfrom = false;
+
+ switch (kind)
+ {
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ copyfrom = true;
+ /* Fallthrough. */
+
+ case GOMP_MAP_TO_PSET:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_DELETE:
+ case GOMP_MAP_RELEASE:
+ {
+ struct splay_tree_key_s cur_node;
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start
+ + (kind == GOMP_MAP_POINTER
+ ? sizeof (void *) : sizes[i]);
+ splay_tree_key n
+ = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ continue;
+
+ if (n->refcount == REFCOUNT_INFINITY)
+ {
+ n->refcount = 1;
+ n->virtual_refcount = 0;
+ }
+
+ if (finalize)
+ {
+ n->refcount -= n->virtual_refcount;
+ n->virtual_refcount = 0;
+ }
+
+ if (n->virtual_refcount > 0)
+ {
+ n->refcount--;
+ n->virtual_refcount--;
+ }
+ else if (n->refcount > 0)
+ n->refcount--;
+
+ if (copyfrom)
+ gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start,
+ (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start),
+ cur_node.host_end - cur_node.host_start);
+
+ if (n->refcount == 0)
+ gomp_remove_var_async (acc_dev, n, aq);
+ }
+ break;
+ default:
+ gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
+ kind);
+ }
+ }
+
+ gomp_mutex_unlock (&acc_dev->lock);
+}
+
void
GOACC_enter_exit_data (int flags_m, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds,
@@ -707,98 +825,13 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum,
va_end (ap);
}
- /* In c, non-pointers and arrays are represented by a single data clause.
- Dynamically allocated arrays and subarrays are represented by a data
- clause followed by an internal GOMP_MAP_POINTER.
-
- In fortran, scalars and not allocated arrays are represented by a
- single data clause. Allocated arrays and subarrays have three mappings:
- 1) the original data clause, 2) a PSET 3) a pointer to the array data.
- */
+ goacc_aq aq = get_goacc_asyncqueue (async);
if (data_enter)
- {
- for (i = 0; i < mapnum; i++)
- {
- unsigned char kind = kinds[i] & 0xff;
-
- /* Scan for pointers and PSETs. */
- int pointer = find_pointer (i, mapnum, kinds);
-
- if (!pointer)
- {
- switch (kind)
- {
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FORCE_ALLOC:
- acc_create_async (hostaddrs[i], sizes[i], async);
- break;
- case GOMP_MAP_TO:
- case GOMP_MAP_FORCE_TO:
- acc_copyin_async (hostaddrs[i], sizes[i], async);
- break;
- default:
- gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
- kind);
- break;
- }
- }
- else
- {
- gomp_acc_insert_pointer (pointer, &hostaddrs[i],
- &sizes[i], &kinds[i], async);
- /* Increment 'i' by two because OpenACC requires fortran
- arrays to be contiguous, so each PSET is associated with
- one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
- one MAP_POINTER. */
- i += pointer - 1;
- }
- }
- }
+ goacc_enter_data_internal (acc_dev, mapnum, hostaddrs, sizes, kinds, aq);
else
- for (i = 0; i < mapnum; ++i)
- {
- unsigned char kind = kinds[i] & 0xff;
-
- int pointer = find_pointer (i, mapnum, kinds);
-
- if (!pointer)
- {
- switch (kind)
- {
- case GOMP_MAP_RELEASE:
- case GOMP_MAP_DELETE:
- if (acc_is_present (hostaddrs[i], sizes[i]))
- {
- if (finalize)
- acc_delete_finalize_async (hostaddrs[i], sizes[i], async);
- else
- acc_delete_async (hostaddrs[i], sizes[i], async);
- }
- break;
- case GOMP_MAP_FROM:
- case GOMP_MAP_FORCE_FROM:
- if (finalize)
- acc_copyout_finalize_async (hostaddrs[i], sizes[i], async);
- else
- acc_copyout_async (hostaddrs[i], sizes[i], async);
- break;
- default:
- gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
- kind);
- break;
- }
- }
- else
- {
- bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
- || kind == GOMP_MAP_FROM);
- gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
- finalize, pointer);
- /* See the above comment. */
- i += pointer - 1;
- }
- }
+ goacc_exit_data_internal (acc_dev, mapnum, hostaddrs, sizes, kinds,
+ finalize, aq);
out_prof:
if (profiling_p)
diff --git a/libgomp/target.c b/libgomp/target.c
index 84d6daa76ca..6ef4156683f 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -536,8 +536,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
- tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
+ tgt->refcount = (pragma_kind == GOMP_MAP_VARS_ENTER_DATA
+ || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
+ tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
if (mapnum == 0)
@@ -937,7 +939,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
- k->dynamic_refcount = 0;
+ k->virtual_refcount = 0;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1075,8 +1077,20 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
- if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
- {
+ if ((pragma_kind == GOMP_MAP_VARS_ENTER_DATA
+ || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA)
+ && tgt->refcount == 0)
+ {
+ /* If we're about to discard a target_mem_desc with no "structural"
+ references (tgt->refcount == 0), any splay keys linked in the tgt's
+ list must have their virtual refcount incremented to represent that
+ "lost" reference in order to implement the semantics of the OpenACC
+ "present increment" operation properly. */
+ if (pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA)
+ for (i = 0; i < tgt->list_count; i++)
+ if (tgt->list[i].key)
+ tgt->list[i].key->virtual_refcount++;
+
free (tgt);
tgt = NULL;
}
@@ -1116,32 +1130,65 @@ gomp_unmap_tgt (struct target_mem_desc *tgt)
free (tgt);
}
-attribute_hidden bool
-gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
+static bool
+gomp_unref_tgt (void *ptr)
{
bool is_tgt_unmapped = false;
- splay_tree_remove (&devicep->mem_map, k);
- if (k->link_key)
- splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
- if (k->tgt->refcount > 1)
- k->tgt->refcount--;
+
+ struct target_mem_desc *tgt = (struct target_mem_desc *) ptr;
+
+ assert (tgt->refcount != REFCOUNT_INFINITY);
+
+ if (tgt->refcount > 1)
+ tgt->refcount--;
else
{
+ gomp_unmap_tgt (tgt);
is_tgt_unmapped = true;
- gomp_unmap_tgt (k->tgt);
}
+
return is_tgt_unmapped;
}
static void
-gomp_unref_tgt (void *ptr)
+gomp_unref_tgt_void (void *ptr)
{
- struct target_mem_desc *tgt = (struct target_mem_desc *) ptr;
+ (void) gomp_unref_tgt (ptr);
+}
- if (tgt->refcount > 1)
- tgt->refcount--;
+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;
+ splay_tree_remove (&devicep->mem_map, k);
+ if (k->link_key)
+ splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
+ if (aq)
+ devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
+ (void *) k->tgt);
else
- gomp_unmap_tgt (tgt);
+ is_tgt_unmapped = gomp_unref_tgt ((void *) k->tgt);
+ return is_tgt_unmapped;
+}
+
+attribute_hidden bool
+gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
+{
+ return gomp_remove_var_internal (devicep, k, NULL);
+}
+
+/* Remove a variable asynchronously. This actually removes the variable
+ mapping immediately, but retains the linked target_mem_desc until the
+ asynchronous operation has completed (as it may still refer to target
+ memory). The device lock must be held before entry, and remains locked on
+ exit. */
+
+attribute_hidden void
+gomp_remove_var_async (struct gomp_device_descr *devicep, splay_tree_key k,
+ struct goacc_asyncqueue *aq)
+{
+ (void) gomp_remove_var_internal (devicep, k, aq);
}
/* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
@@ -1177,7 +1224,14 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
continue;
bool do_unmap = false;
- if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ if (k->tgt == tgt
+ && k->virtual_refcount > 0
+ && k->refcount != REFCOUNT_INFINITY)
+ {
+ k->virtual_refcount--;
+ k->refcount--;
+ }
+ else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
else if (k->refcount == 1)
{
@@ -1197,7 +1251,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
}
if (aq)
- devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt,
+ devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
(void *) tgt);
else
gomp_unref_tgt ((void *) tgt);
@@ -1334,6 +1388,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
+ k->virtual_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -1366,6 +1421,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_var->start;
k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
+ k->virtual_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -1600,22 +1656,6 @@ gomp_unload_device (struct gomp_device_descr *devicep)
}
}
-/* Free address mapping tables. MM must be locked on entry, and remains locked
- on return. */
-
-attribute_hidden void
-gomp_free_memmap (struct splay_tree_s *mem_map)
-{
- while (mem_map->root)
- {
- struct target_mem_desc *tgt = mem_map->root->key.tgt;
-
- splay_tree_remove (mem_map, &mem_map->root->key);
- free (tgt->array);
- free (tgt);
- }
-}
-
/* Host fallback for GOMP_target{,_ext} routines. */
static void
@@ -2100,6 +2140,7 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
if (k->link_key)
splay_tree_insert (&devicep->mem_map,
(splay_tree_node) k->link_key);
+ assert (k->tgt->refcount != REFCOUNT_INFINITY);
if (k->tgt->refcount > 1)
k->tgt->refcount--;
else
@@ -2636,6 +2677,8 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
k->tgt = tgt;
k->tgt_offset = (uintptr_t) device_ptr + device_offset;
k->refcount = REFCOUNT_INFINITY;
+ k->virtual_refcount = 0;
+ k->link_key = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
@@ -2906,7 +2949,6 @@ gomp_target_init (void)
current_device.type = current_device.get_type_func ();
current_device.mem_map.root = NULL;
current_device.state = GOMP_DEVICE_UNINITIALIZED;
- current_device.openacc.data_environ = NULL;
for (i = 0; i < new_num_devices; i++)
{
current_device.target_id = i;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/unmap-infinity-2.c b/libgomp/testsuite/libgomp.c-c++-common/unmap-infinity-2.c
new file mode 100644
index 00000000000..3931c5aba25
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/unmap-infinity-2.c
@@ -0,0 +1,19 @@
+int foo[16];
+#pragma omp declare target (foo)
+
+__attribute__((used)) void bar (void)
+{
+ #pragma omp target parallel for
+ for (int i = 0; i < 16; i++)
+ foo[i] = i;
+}
+
+int
+main (int argc, char *argv[])
+{
+ int *foo_copy = foo;
+ /* Try to trigger the unmapping of a REFCOUNT_INFINITY target block. This
+ does nothing at the time of writing. */
+ #pragma omp target exit data map(delete: foo_copy[0:16])
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c
new file mode 100644
index 00000000000..872f0c1de5c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int foo[16];
+#pragma acc declare device_resident(foo)
+
+int
+main (int argc, char *argv[])
+{
+ acc_init (acc_device_default);
+ acc_unmap_data ((void *) foo);
+/* { dg-output "libgomp: cannot unmap target block" } */
+ return 0;
+}
+
+/* { dg-shouldfail "" } */
