From: Julian Brown <julian@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: Thomas Schwinge <thomas@codesourcery.com>,
Jakub Jelinek <jakub@redhat.com>,
Tobias Burnus <tobias@codesourcery.com>,
<Catherine_Moore@mentor.com>, <fortran@gcc.gnu.org>
Subject: [PATCH 02/13] OpenACC reference count overhaul
Date: Wed, 18 Dec 2019 06:03:00 -0000 [thread overview]
Message-ID: <491e3ca360313930f8f2f5686ffd386cf2fad04e.1576648001.git.julian@codesourcery.com> (raw)
In-Reply-To: <cover.1576648001.git.julian@codesourcery.com>
This is a rebased version of the reference-count overhaul patch last
posted here:
https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02235.html
This version omits parts of the above patch already committed upstream and
merges some recent REFCOUNT_INFINITY changes. This patch causes the newish
PR92843 test to fail, though IMO that test relies on behaviour arising
from a rather nuanced reading of the spec. Hopefully we can resolve that
problem as a follow-up.
Tested alongside other patches in this series with offloading to
NVPTX. OK?
Julian
2019-11-22 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* libgomp.h (struct splay_tree_key_s): Substitute dynamic_refcount
field for virtual_refcount.
(enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA.
(gomp_free_memmap): Remove prototype.
* oacc-init.c (acc_shutdown_1): Iteratively call gomp_remove_var
instead of calling gomp_free_memmap.
* oacc-mem.c (acc_unmap_data): Open code instead of forcing
target_mem_desc's to_free NULL then calling gomp_unmap_vars. Handle
REFCOUNT_INFINITY on target blocks.
(present_create_copy): Use virtual_refcount instead of
dynamic_refcount. Re-do lookup for target pointer return value.
(delete_copyout): Update for virtual_refcount semantics.
(gomp_acc_insert_pointer, gomp_acc_remove_pointer, find_pointer):
Remove functions.
(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_unmap_vars_internal): Update for virtual_refcount semantics.
(gomp_load_image_to_device, omp_target_associate_ptr): Zero-initialise
virtual_refcount field instead of dynamic_refcount.
(gomp_free_memmap): Remove function.
* testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c: New test.
* testsuite/libgomp.c-c++-common/unmap-infinity-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c:
Remove PR92848 TODOs.
* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Add XFAIL.
---
libgomp/libgomp.h | 9 +-
libgomp/oacc-init.c | 10 +-
libgomp/oacc-mem.c | 399 +++++++-----------
libgomp/target.c | 53 +--
.../libgomp.c-c++-common/unmap-infinity-2.c | 19 +
.../libgomp.oacc-c-c++-common/pr92843-1.c | 1 +
.../subset-subarray-mappings-1-r-p.c | 16 -
.../unmap-infinity-1.c | 17 +
8 files changed, 228 insertions(+), 296 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/unmap-infinity-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 0f1f11284d5..865b9df2444 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1007,8 +1007,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;
struct splay_tree_aux *aux;
};
@@ -1139,6 +1142,7 @@ 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
@@ -1169,7 +1173,6 @@ 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,
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index a444c604d59..dd88b58a379 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -370,7 +370,15 @@ 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;
+ if (k->aux)
+ k->aux->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 196b7e2a520..2a0e7236b92 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -406,7 +406,7 @@ acc_map_data (void *h, void *d, size_t s)
&kinds, true, GOMP_MAP_VARS_OPENACC);
splay_tree_key n = tgt->list[0].key;
assert (n->refcount == 1);
- assert (n->dynamic_refcount == 0);
+ assert (n->virtual_refcount == 0);
/* Special reference counting behavior. */
n->refcount = REFCOUNT_INFINITY;
@@ -434,12 +434,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)
{
@@ -447,7 +444,7 @@ acc_unmap_data (void *h)
gomp_fatal ("%p is not a mapped block", (void *)h);
}
- host_size = n->host_end - n->host_start;
+ size_t host_size = n->host_end - n->host_start;
if (n->host_start != (uintptr_t) h)
{
@@ -456,7 +453,7 @@ acc_unmap_data (void *h)
(void *) n->host_start, (int) host_size, (void *) h);
}
/* TODO This currently doesn't catch 'REFCOUNT_INFINITY' usage different from
- 'acc_map_data'. Maybe 'dynamic_refcount' can be used for disambiguating
+ 'acc_map_data'. Maybe 'virtual_refcount' can be used for disambiguating
the different 'REFCOUNT_INFINITY' cases, or simply separate
'REFCOUNT_INFINITY' values per different usage ('REFCOUNT_ACC_MAP_DATA'
etc.)? */
@@ -468,24 +465,25 @@ acc_unmap_data (void *h)
(void *) h, (int) host_size);
}
- /* 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)
{
- /* 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;
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("cannot unmap target block");
+ }
+ else if (tgt->refcount > 1)
+ tgt->refcount--;
+ else
+ {
+ free (tgt->array);
+ free (tgt);
}
gomp_mutex_unlock (&acc_dev->lock);
- gomp_unmap_vars (t, true);
-
if (profiling_p)
{
thr->prof_info = NULL;
@@ -545,8 +543,10 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
assert (n->refcount != REFCOUNT_LINK);
if (n->refcount != REFCOUNT_INFINITY)
- n->refcount++;
- n->dynamic_refcount++;
+ {
+ n->refcount++;
+ n->virtual_refcount++;
+ }
gomp_mutex_unlock (&acc_dev->lock);
}
@@ -557,7 +557,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;
@@ -571,14 +570,16 @@ 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);
- n = tgt->list[0].key;
- assert (n->refcount == 1);
- assert (n->dynamic_refcount == 0);
- n->dynamic_refcount++;
+ gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s, &kinds,
+ true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
- d = tgt->to_free;
+ gomp_mutex_lock (&acc_dev->lock);
+ 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);
}
if (profiling_p)
@@ -696,26 +697,21 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
(void *) h, (int) s, (void *) n->host_start, (int) host_size);
}
- assert (n->refcount != REFCOUNT_LINK);
- if (n->refcount != REFCOUNT_INFINITY
- && n->refcount < n->dynamic_refcount)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
- }
-
if (f & FLAG_FINALIZE)
{
if (n->refcount != REFCOUNT_INFINITY)
- 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)
{
if (n->refcount != REFCOUNT_INFINITY)
n->refcount--;
- n->dynamic_refcount--;
+ n->virtual_refcount--;
}
+ else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ n->refcount--;
if (n->refcount == 0)
{
@@ -870,154 +866,138 @@ acc_update_self_async (void *h, size_t s, int async)
update_dev_host (0, h, s, async);
}
+/* 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. */
-/* OpenACC 'enter data', 'exit data': 'GOACC_enter_exit_data' and its helper
- functions. */
-
-/* Special handling for 'GOMP_MAP_POINTER', 'GOMP_MAP_TO_PSET'.
-
- Only the first mapping is considered in reference counting; the following
- ones implicitly follow suit. */
-
-static void
-goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
- void *kinds, int async)
+static int
+find_group_last (int pos, size_t mapnum, unsigned short *kinds)
{
- struct target_mem_desc *tgt;
- struct goacc_thread *thr = goacc_thread ();
- struct gomp_device_descr *acc_dev = thr->dev;
-
- if (*hostaddrs == NULL)
- return;
+ unsigned char kind0 = kinds[pos] & 0xff;
+ int first_pos = pos, last_pos = pos;
- if (acc_is_present (*hostaddrs, *sizes))
+ if (kind0 == GOMP_MAP_TO_PSET)
{
- splay_tree_key n;
- gomp_mutex_lock (&acc_dev->lock);
- n = lookup_host (acc_dev, *hostaddrs, *sizes);
- assert (n->refcount != REFCOUNT_INFINITY
- && n->refcount != REFCOUNT_LINK);
- 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");
+ 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;
}
- 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);
- splay_tree_key n = tgt->list[0].key;
- assert (n->refcount == 1);
- assert (n->dynamic_refcount == 0);
- n->dynamic_refcount++;
- gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
+ return last_pos;
}
+/* 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_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
- int finalize, int mapnum)
+goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
+ void **hostaddrs, size_t *sizes,
+ unsigned short *kinds, goacc_aq aq)
{
- 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)
+ for (size_t i = 0; i < mapnum; i++)
{
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("%p is not a mapped block", (void *)h);
- }
-
- gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
+ int group_last = find_group_last (i, mapnum, kinds);
- t = n->tgt;
+ 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);
- assert (n->refcount != REFCOUNT_INFINITY
- && n->refcount != REFCOUNT_LINK);
- if (n->refcount < n->dynamic_refcount)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
+ i = group_last;
}
+}
- if (finalize)
- {
- n->refcount -= n->dynamic_refcount;
- n->dynamic_refcount = 0;
- }
- else if (n->dynamic_refcount)
- {
- n->refcount--;
- n->dynamic_refcount--;
- }
+/* Unmap variables for OpenACC "exit data", with optional finalization
+ (affecting all mappings in this operation). */
- gomp_mutex_unlock (&acc_dev->lock);
+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);
- if (n->refcount == 0)
+ for (size_t i = 0; i < mapnum; ++i)
{
- /* 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;
- }
+ unsigned char kind = kinds[i] & 0xff;
+ bool copyfrom = false;
- /* If running synchronously, unmap immediately. */
- if (async < acc_async_noval)
- gomp_unmap_vars (t, true);
- else
+ switch (kind)
{
- goacc_aq aq = get_goacc_asyncqueue (async);
- gomp_unmap_vars_async (t, true, aq);
+ 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 (finalize)
+ {
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount -= n->virtual_refcount;
+ n->virtual_refcount = 0;
+ }
+
+ if (n->virtual_refcount > 0)
+ {
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount--;
+ n->virtual_refcount--;
+ }
+ else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ n->refcount--;
+
+ if (copyfrom
+ && (kind != GOMP_MAP_FROM || n->refcount == 0))
+ 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);
-
- gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
-}
-
-/* Return the number of mappings associated with 'GOMP_MAP_TO_PSET' or
- 'GOMP_MAP_POINTER'. */
-
-static int
-find_pointer (int pos, size_t mapnum, unsigned short *kinds)
-{
- if (pos + 1 >= mapnum)
- return 0;
-
- unsigned char kind = kinds[pos+1] & 0xff;
-
- if (kind == GOMP_MAP_TO_PSET)
- return 3;
- else if (kind == GOMP_MAP_POINTER)
- return 2;
-
- return 0;
}
void
@@ -1147,98 +1127,13 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
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
- {
- goacc_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);
- goacc_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 97c2b5c5e4d..23f9e1618ca 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)
@@ -939,7 +941,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;
@@ -1077,8 +1079,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;
}
@@ -1216,7 +1230,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)
{
@@ -1373,7 +1394,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->dynamic_refcount = 0;
+ k->virtual_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -1406,7 +1427,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->dynamic_refcount = 0;
+ k->virtual_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -1641,22 +1662,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
@@ -2668,7 +2673,7 @@ 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->dynamic_refcount = 0;
+ k->virtual_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
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/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
index db5b35b08d9..f16c46a37bf 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -1,6 +1,7 @@
/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
reference count. */
+/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
index 9b5d83c66dd..907b8587773 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c
@@ -156,20 +156,16 @@ f1 (void)
assert (acc_is_present (&myblock[i], SUBSET));
assert (acc_is_present (myblock, SIZE));
-#if 0 //TODO PR92848
if (last)
cb_ev_free_expected = true;
-#endif
#if OPENACC_RUNTIME
acc_delete (&myblock[i], SUBSET);
#else
# pragma acc exit data delete (myblock[i:SUBSET])
#endif
-#if 0 //TODO PR92848
assert (!cb_ev_free_expected);
if (last)
assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
-#endif
assert (acc_is_present (&myblock[i], SUBSET) != last);
assert (acc_is_present (myblock, SIZE) != last);
}
@@ -331,9 +327,7 @@ f3 ()
assert (acc_is_present (h, SIZE));
assert (acc_is_present (&h[2], SIZE - 2));
-#if 0 //TODO PR92848
cb_ev_free_expected = true;
-#endif
#if OPENACC_RUNTIME
acc_delete (h, SIZE);
#else
@@ -343,10 +337,8 @@ f3 ()
# pragma acc exit data delete (h)
# endif
#endif
-#if 0 //TODO PR92848
assert (!cb_ev_free_expected);
assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
-#endif
assert (!acc_is_present (h, SIZE));
assert (!acc_is_present (&h[2], SIZE - 2));
@@ -401,19 +393,15 @@ f_lib_22 (void)
memset (h, c1, SIZE);
/* Now 'copyout' not the whole but only a "subset" subarray, missing one
SUBSET at the beginning, and half a SUBSET at the end... */
-#if 0 //TODO PR92848
cb_ev_free_expected = true;
-#endif
#if OPENACC_RUNTIME
acc_copyout (h + SUBSET, SIZE - SUBSET - SUBSET / 2);
#else
# pragma acc exit data copyout (h[SUBSET:SIZE - SUBSET - SUBSET / 2])
#endif
-#if 0 //TODO PR92848
/* ..., yet, expect the device memory object to be 'free'd... */
assert (!cb_ev_free_expected);
assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
-#endif
/* ..., and the mapping to be removed... */
assert (!acc_is_present (h, SIZE));
assert (!acc_is_present (&h[SUBSET], SIZE - SUBSET - SUBSET / 2));
@@ -474,19 +462,15 @@ f_lib_30 (void)
assert (aligned_address (cb_ev_alloc_device_ptr) == d);
/* We 'delete' not the whole but only a "subset" subarray... */
-#if 0 //TODO PR92848
cb_ev_free_expected = true;
-#endif
#if OPENACC_RUNTIME
acc_delete (h, SIZE - SUBSET);
#else
# pragma acc exit data delete (h[0:SIZE - SUBSET])
#endif
-#if 0 //TODO PR92848
/* ..., yet, expect the device memory object to be 'free'd... */
assert (!cb_ev_free_expected);
assert (cb_ev_free_device_ptr == cb_ev_alloc_device_ptr);
-#endif
/* ..., and the mapping to be removed. */
assert (!acc_is_present (h, SIZE));
assert (!acc_is_present (h, SIZE - SUBSET));
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 "" } */
--
2.23.0
next prev parent reply other threads:[~2019-12-18 6:03 UTC|newest]
Thread overview: 81+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-11-10 17:11 [PATCH 0/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-10 17:11 ` [PATCH 2/3] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2018-12-18 14:16 ` Julian Brown
2018-12-18 14:50 ` Jakub Jelinek
2018-11-10 17:11 ` [PATCH 1/3] Host-to-device transfer coalescing & magic offset value self-documentation Julian Brown
2018-12-21 10:56 ` libgomp/target.c magic constants self-documentation Thomas Schwinge
2019-05-29 14:48 ` Thomas Schwinge
2018-11-10 17:12 ` [PATCH 3/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-11 17:04 ` Bernhard Reutner-Fischer
2018-11-30 11:41 ` [PATCH] " Julian Brown
2018-12-03 17:03 ` Julian Brown
2018-12-07 13:50 ` Jakub Jelinek
2018-12-10 19:42 ` Julian Brown
2018-12-13 10:57 ` Jakub Jelinek
2018-12-14 19:00 ` Julian Brown
2018-12-18 12:25 ` Jakub Jelinek
2018-12-22 13:37 ` Thomas Schwinge
2019-10-18 17:20 ` Thomas Schwinge
2019-11-06 18:44 ` Julian Brown
2019-11-22 23:54 ` Julian Brown
2019-11-25 10:53 ` Tobias Burnus
2019-11-26 2:54 ` Julian Brown
2019-12-17 12:16 ` Thomas Schwinge
2019-12-17 17:28 ` [WIP] OpenACC 'acc_attach*', 'acc_detach*' runtime library routines (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
2019-12-18 6:03 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Julian Brown
2019-12-18 6:03 ` [PATCH 01/13] Use aux struct in libgomp for infrequently-used/API-specific data Julian Brown
2019-12-18 6:03 ` Julian Brown [this message]
2020-05-19 15:42 ` [PATCH 02/13] OpenACC reference count overhaul Thomas Schwinge
2020-06-04 18:13 ` [OpenACC] Use 'tgt' returned from 'gomp_map_vars' (was: [PATCH 02/13] OpenACC reference count overhaul) Thomas Schwinge
2020-05-19 15:49 ` [PATCH 02/13] OpenACC reference count overhaul Thomas Schwinge
2020-05-19 15:58 ` Thomas Schwinge
2020-06-25 11:03 ` Thomas Schwinge
2020-07-03 15:29 ` Thomas Schwinge
2019-12-18 6:03 ` [PATCH 03/13] OpenACC reference count consistency checking Julian Brown
2019-12-18 6:04 ` [PATCH 08/13] OpenACC 2.6 deep copy: middle-end parts Julian Brown
2019-12-21 21:51 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts Julian Brown
2019-12-24 5:05 ` Thomas Schwinge
2019-12-26 19:04 ` Jason Merrill
2021-06-10 11:03 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 06/13] OpenACC 2.6 deep copy: attach/detach API routines Julian Brown
2019-12-18 6:04 ` [PATCH 05/13] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2019-12-18 6:04 ` [PATCH 04/13] Use gomp_map_val for OpenACC host-to-device address translation Julian Brown
2019-12-18 6:05 ` [PATCH 13/13] Fortran polymorphic class-type support for OpenACC Julian Brown
2019-12-18 6:05 ` [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests Julian Brown
2020-06-04 18:43 ` Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7, 8}.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests) Thomas Schwinge
2023-10-31 14:00 ` Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' " Thomas Schwinge
2019-12-18 6:05 ` [PATCH 12/13] OpenACC 2.6 deep copy: Fortran execution tests Julian Brown
2019-12-18 6:05 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Julian Brown
2019-12-21 23:37 ` Thomas Schwinge
2020-01-03 12:26 ` Julian Brown
2020-05-20 9:37 ` Thomas Schwinge
2020-06-05 16:23 ` [OpenACC 'exit data'] Simplify 'GOMP_MAP_STRUCT' handling (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-05 16:36 ` [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings " Thomas Schwinge
2020-05-20 14:52 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-05-20 19:11 ` Julian Brown
2020-06-04 18:35 ` [OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-04 18:53 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-06-05 10:39 ` Thomas Schwinge
2020-06-05 20:28 ` Julian Brown
2020-06-05 11:17 ` Thomas Schwinge
2020-06-05 20:31 ` Julian Brown
2020-06-09 10:41 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-09 12:23 ` Julian Brown
2020-06-18 18:21 ` Julian Brown
2020-07-16 8:35 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting Thomas Schwinge
2020-06-26 9:20 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-07-16 9:35 ` Thomas Schwinge
2020-07-16 21:21 ` Julian Brown
2020-07-17 9:12 ` Thomas Schwinge
2020-06-30 15:58 ` Thomas Schwinge
2019-12-18 7:20 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Julian Brown
2019-12-18 23:30 ` Tobias Burnus
2019-12-20 12:25 ` [committed] Improve is-coindexed check for OpenACC/OpenMP (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-20 13:25 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Tobias Burnus
2019-12-20 10:08 ` [patch,committed] Fix testsuite-fallout of OpenACC deep-copy patch (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-18 18:24 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Thomas Schwinge
2019-12-20 1:21 ` Julian Brown
2019-12-20 14:36 ` OpenACC regression and development pace Thomas Koenig
2020-06-04 18:07 ` [OpenACC] XFAIL behavior of over-eager 'finalize' clause (was: [PATCH 00/13] OpenACC 2.6 manual deep copy support) Thomas Schwinge
2019-12-17 16:53 ` In 'libgomp/target.c', 'struct splay_tree_key_s', use 'struct splay_tree_aux' for infrequently-used or API-specific data (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=491e3ca360313930f8f2f5686ffd386cf2fad04e.1576648001.git.julian@codesourcery.com \
--to=julian@codesourcery.com \
--cc=Catherine_Moore@mentor.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=thomas@codesourcery.com \
--cc=tobias@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).