diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index ca42e0de64041621079678b9d82f5b2101d76591..7b52ce7d5c23306141a29b45bac57788225fcfb1 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1016,11 +1016,8 @@ struct splay_tree_key_s {
uintptr_t tgt_offset;
/* Reference count. */
uintptr_t 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;
+ /* Dynamic reference count. */
+ uintptr_t dynamic_refcount;
struct splay_tree_aux *aux;
};
@@ -1153,7 +1150,6 @@ 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
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 34f519a2045292736fca82a90e9368c84e453a85..855cad84391192ba9fe65d019e87a58d9ce0665c 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -409,7 +409,7 @@ acc_map_data (void *h, void *d, size_t s)
splay_tree_key n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
- assert (n->virtual_refcount == 0);
+ assert (n->dynamic_refcount == 0);
/* Special reference counting behavior. */
n->refcount = REFCOUNT_INFINITY;
@@ -456,7 +456,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 'virtual_refcount' can be used for disambiguating
+ 'acc_map_data'. Maybe 'dynamic_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.)? */
@@ -520,10 +520,8 @@ goacc_map_var_existing (struct gomp_device_descr *acc_dev, void *hostaddr,
assert (n->refcount != REFCOUNT_LINK);
if (n->refcount != REFCOUNT_INFINITY)
- {
- n->refcount++;
- n->virtual_refcount++;
- }
+ n->refcount++;
+ n->dynamic_refcount++;
return d;
}
@@ -574,13 +572,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ kinds, true, GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
- assert (n->virtual_refcount == 0);
+ assert (n->dynamic_refcount == 0);
+ n->dynamic_refcount++;
d = (void *) tgt->tgt_start;
}
@@ -676,24 +675,30 @@ goacc_exit_datum_1 (struct gomp_device_descr *acc_dev, void *h, size_t s,
(void *) h, (int) s, (void *) n->host_start, (int) host_size);
}
- bool finalize = (kind == GOMP_MAP_DELETE
- || kind == GOMP_MAP_FORCE_FROM);
+ bool finalize = (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FORCE_DETACH);
+
+ 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 (finalize)
{
if (n->refcount != REFCOUNT_INFINITY)
- n->refcount -= n->virtual_refcount;
- n->virtual_refcount = 0;
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
}
-
- if (n->virtual_refcount > 0)
+ else if (n->dynamic_refcount)
{
if (n->refcount != REFCOUNT_INFINITY)
n->refcount--;
- n->virtual_refcount--;
+ n->dynamic_refcount--;
}
- else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
- n->refcount--;
if (n->refcount == 0)
{
@@ -1068,18 +1073,144 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
void **hostaddrs, size_t *sizes,
unsigned short *kinds, goacc_aq aq)
{
+ gomp_mutex_lock (&acc_dev->lock);
+
for (size_t i = 0; i < mapnum; i++)
{
- int group_last = find_group_last (i, mapnum, sizes, kinds);
+ splay_tree_key n;
+ size_t group_last = find_group_last (i, mapnum, sizes, kinds);
+ bool struct_p = false;
+ size_t size, groupnum = (group_last - i) + 1;
+
+ switch (kinds[i] & 0xff)
+ {
+ case GOMP_MAP_STRUCT:
+ {
+ size = (uintptr_t) hostaddrs[group_last] + sizes[group_last]
+ - (uintptr_t) hostaddrs[i];
+ struct_p = true;
+ }
+ break;
+
+ case GOMP_MAP_ATTACH:
+ size = sizeof (void *);
+ break;
+
+ default:
+ size = sizes[i];
+ }
- 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);
+ n = lookup_host (acc_dev, hostaddrs[i], size);
+
+ if (n && struct_p)
+ {
+ for (size_t j = i + 1; j <= group_last; j++)
+ {
+ struct splay_tree_key_s cur_node;
+ cur_node.host_start = (uintptr_t) hostaddrs[j];
+ cur_node.host_end = cur_node.host_start + sizes[j];
+ splay_tree_key n2
+ = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+ if (!n2
+ || n2->tgt != n->tgt
+ || n2->host_start - n->host_start
+ != n2->tgt_offset - n->tgt_offset)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("Trying to map into device [%p..%p) structure "
+ "element when other mapped elements from the "
+ "same structure weren't mapped together with "
+ "it", (void *) cur_node.host_start,
+ (void *) cur_node.host_end);
+ }
+ }
+ /* This is a special case because we must increment the refcount by
+ the number of mapped struct elements, rather than by one. */
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount += groupnum - 1;
+ n->dynamic_refcount += groupnum - 1;
+ }
+ else if (n && groupnum == 1)
+ {
+ void *h = hostaddrs[i];
+ size_t s = sizes[i];
+
+ /* A standalone attach clause. */
+ if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
+ gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
+ (uintptr_t) h, s, NULL);
+
+ goacc_map_var_existing (acc_dev, h, s, n);
+ }
+ else if (n && groupnum > 1)
+ {
+ assert (n->refcount != REFCOUNT_INFINITY
+ && n->refcount != REFCOUNT_LINK);
+
+ for (size_t j = i + 1; j <= group_last; j++)
+ if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH)
+ {
+ splay_tree_key m
+ = lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
+ gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
+ (uintptr_t) hostaddrs[j], sizes[j], NULL);
+ }
+
+ bool processed = false;
+
+ struct target_mem_desc *tgt = n->tgt;
+ for (size_t j = 0; j < tgt->list_count; j++)
+ if (tgt->list[j].key == n)
+ {
+ /* We are processing a group of mappings (e.g.
+ [GOMP_MAP_TO, GOMP_MAP_TO_PSET, GOMP_MAP_POINTER]).
+ Find the right group in the target_mem_desc's variable
+ list, and increment the refcounts for each item in that
+ group. */
+ for (size_t k = 0; k < groupnum; k++)
+ if (j + k < tgt->list_count && tgt->list[j + k].key)
+ {
+ tgt->list[j + k].key->refcount++;
+ tgt->list[j + k].key->dynamic_refcount++;
+ }
+ processed = true;
+ break;
+ }
+
+ if (!processed)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("dynamic refcount incrementing failed for "
+ "pointer/pset");
+ }
+ }
+ else if (hostaddrs[i])
+ {
+ /* The data is not mapped already. Map it now, unless the first
+ member in the group has a NULL pointer (e.g. a non-present
+ optional parameter). */
+ gomp_mutex_unlock (&acc_dev->lock);
+
+ struct target_mem_desc *tgt
+ = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
+ &sizes[i], &kinds[i], true,
+ GOMP_MAP_VARS_ENTER_DATA);
+ assert (tgt);
+
+ gomp_mutex_lock (&acc_dev->lock);
+
+ for (size_t j = 0; j < tgt->list_count; j++)
+ {
+ n = tgt->list[j].key;
+ if (n)
+ n->dynamic_refcount++;
+ }
+ }
i = group_last;
}
+
+ gomp_mutex_unlock (&acc_dev->lock);
}
/* Unmap variables for OpenACC "exit data". */
@@ -1128,21 +1259,11 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
for (size_t i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
- bool copyfrom = false;
- bool finalize = false;
-
- if (kind == GOMP_MAP_FORCE_FROM
- || kind == GOMP_MAP_DELETE
- || kind == GOMP_MAP_FORCE_DETACH)
- finalize = true;
switch (kind)
{
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
- copyfrom = true;
- /* Fallthrough. */
-
case GOMP_MAP_TO_PSET:
case GOMP_MAP_POINTER:
case GOMP_MAP_DELETE:
@@ -1166,54 +1287,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
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 (n->refcount == 0)
- {
- if (copyfrom)
- {
- void *d = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start - n->host_start);
- gomp_copy_dev2host (acc_dev, aq,
- (void *) cur_node.host_start, d,
- cur_node.host_end - cur_node.host_start);
- }
-
- if (aq)
- /* TODO We can't do the 'is_tgt_unmapped' checking -- see the
- 'gomp_unref_tgt' comment in
- <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
- PR92881. */
- gomp_remove_var_async (acc_dev, n, aq);
- else
- {
- size_t num_mappings = 0;
- /* If the target_mem_desc represents a single data mapping,
- we can check that it is freed when this splay tree key's
- refcount reaches zero. Otherwise (e.g. for a
- 'GOMP_MAP_STRUCT' mapping with multiple members), fall
- back to skipping the test. */
- for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
- if (n->tgt->list[l_i].key)
- ++num_mappings;
- bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
- assert (is_tgt_unmapped || num_mappings > 1);
- }
- }
+ goacc_exit_datum_1 (acc_dev, hostaddrs[i], size, kind, n, aq);
}
break;
diff --git a/libgomp/target.c b/libgomp/target.c
index d4a4a408b400efde3c089d55305c4dfc1dd64cc5..d6b3572c8d88e649cefe005d7e7831c897397958 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -668,8 +668,7 @@ 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
- || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1;
+ tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -1095,7 +1094,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].copy_from = false;
tgt->list[i].always_copy_from = false;
tgt->list[i].do_detach
- = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ = (pragma_kind != GOMP_MAP_VARS_ENTER_DATA);
n->refcount++;
}
else
@@ -1156,7 +1155,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->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1295,20 +1294,8 @@ 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
- || 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++;
-
+ if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+ {
free (tgt);
tgt = NULL;
}
@@ -1460,14 +1447,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
continue;
bool do_unmap = false;
- 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)
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
else if (k->refcount == 1)
{
@@ -1632,7 +1612,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->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -1666,7 +1646,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_var->start;
k->refcount = is_link_var ? REFCOUNT_LINK : REFCOUNT_INFINITY;
- k->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
@@ -2936,7 +2916,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->virtual_refcount = 0;
+ k->dynamic_refcount = 0;
k->aux = NULL;
array->left = NULL;
array->right = NULL;
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 78fe1402ad464fce5d4ce5a3f48254f0ef3f9314..db5b35b08d9f56b95eb8dae42d07141540f816b5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
@@ -4,7 +4,6 @@
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
-#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
@@ -135,15 +134,7 @@ test_acc_data ()
assert (acc_is_present (h, sizeof h));
assign_array (h, N, c1);
- fprintf (stderr, "CheCKpOInT1\n");
- // { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
acc_copyout_finalize (h, sizeof h);
- //TODO goacc_exit_datum: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- //TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- //TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- //TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- fprintf (stderr, "CheCKpOInT2\n");
- // { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
new file mode 100644
index 0000000000000000000000000000000000000000..4e6d06d48d5c902a3334deae69ae87501b7d6f75
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c
@@ -0,0 +1,31 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ char b;
+};
+
+int main ()
+{
+ struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+
+ assert (acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+
+#pragma acc exit data delete(s.a)
+#pragma acc exit data delete(s.b)
+
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (!acc_is_present (&s.b, sizeof s.b));
+
+ return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
new file mode 100644
index 0000000000000000000000000000000000000000..5539fd8d57f305a2ca41e8f256249159f3016546
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c
@@ -0,0 +1,31 @@
+/* Test dynamic unmapping of separate structure members. */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ char b;
+};
+
+int main ()
+{
+ struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+
+ assert (acc_is_present (&s.a, sizeof s.a));
+ assert (acc_is_present (&s.b, sizeof s.b));
+
+ acc_delete (&s.a, sizeof s.a);
+ acc_delete (&s.b, sizeof s.b);
+
+ assert (!acc_is_present (&s.a, sizeof s.a));
+ assert (!acc_is_present (&s.b, sizeof s.b));
+
+ return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c
new file mode 100644
index 0000000000000000000000000000000000000000..a644ea9c26f389e872c2f576c9f6cabca8bdffb6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c
@@ -0,0 +1,34 @@
+/* Test dynamic mapping of separate structure members. */
+
+#include <assert.h>
+#include <stdio.h>
+#include <openacc.h>
+
+struct s
+{
+ char a;
+ float b;
+};
+
+int main ()
+{
+ struct s s;
+
+#pragma acc enter data create(s.a)
+ assert (acc_is_present (&s.a, sizeof s.a));
+
+ fprintf (stderr, "CheCKpOInT1\n");
+ /* { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } */
+#pragma acc enter data create(s.b)
+ /* { dg-output "(\n|\r\n|\r)libgomp: Trying to map into device \\\[\[0-9a-fA-FxX.\]+\\\) structure element when other mapped elements from the same structure weren't mapped together with it(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+ { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log. */
+ fprintf (stderr, "CheCKpOInT2\n");
+ /* { dg-output "CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } } */
+ assert (acc_is_present (&s.b, sizeof s.b));
+
+ //TODO PR95236
+ assert (acc_is_present (&s, sizeof s));
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
index e9a6510ace826073b43cdbc9c7c89e8db46f7501..09d2ad54e8781edcdb048e8378b4b68f59747fd7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
@@ -13,8 +13,6 @@ main (int argc, char *argv[])
char *block2 = (char *) malloc (SIZE);
char *block3 = (char *) malloc (SIZE);
- /* Doing this twice ensures that we have a non-zero virtual refcount. Make
- sure that works too. */
#ifdef OPENACC_API
acc_copyin (block1, SIZE);
acc_copyin (block1, SIZE);
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
index ed4f10e7a3f88f31625b805451b645362bc2585f..038f04a3c37e46e5fb6c54f30a6e8502a0f29d28 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
@@ -1,8 +1,5 @@
! { dg-do run }
-/* Nullify the 'finalize' clause, which disturbs reference counting. */
+/* Nullify the 'finalize' clause. */
#define finalize
#include "deep-copy-6.f90"
-
-! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
index 5837a4039106045ed43dbf1c26ec09f94e03b29c..6aab6a4a763378473c5c7c3d0a26499acbc469a1 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -40,15 +40,7 @@ program dtype
if (.not. acc_is_present(var%a(5:n - 5))) stop 11
if (.not. acc_is_present(var%b(5:n - 5))) stop 12
if (.not. acc_is_present(var)) stop 13
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_get_device_type() .ne. acc_device_host) then
if (acc_is_present(var%a(5:n - 5))) stop 21
if (acc_is_present(var%b(5:n - 5))) stop 22
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90
new file mode 100644
index 0000000000000000000000000000000000000000..6b17b1dbbc920ffc3f4405dcc30707944102d8a7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90
@@ -0,0 +1,49 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program map_multi
+ use openacc
+ implicit none
+ integer, parameter :: n = 512
+ integer, allocatable :: a(:), b(:), c(:)
+
+ allocate(a(1:n))
+ allocate(b(1:n))
+ allocate(c(1:n))
+
+ !$acc data copy(a, b, c)
+
+ ! These arrays have descriptors, so use multiple mappings. Make sure those
+ ! are matched up properly with the mappings in the enclosing data region.
+ !$acc enter data copyin(a)
+ !$acc enter data copyin(b)
+ !$acc enter data copyin(c)
+
+ !$acc end data
+
+ if (.not.acc_is_present (a)) stop 1
+ if (.not.acc_is_present (b)) stop 2
+ if (.not.acc_is_present (c)) stop 3
+
+ !$acc exit data delete(a)
+
+ if (acc_is_present (a)) stop 4
+ if (.not.acc_is_present (b)) stop 5
+ if (.not.acc_is_present (c)) stop 6
+
+ !$acc exit data delete(b)
+
+ if (acc_is_present (a)) stop 7
+ if (acc_is_present (b)) stop 8
+ if (.not.acc_is_present (c)) stop 9
+
+ !$acc exit data delete(c)
+
+ if (acc_is_present (a)) stop 10
+ if (acc_is_present (b)) stop 11
+ if (acc_is_present (c)) stop 12
+
+ deallocate(a)
+ deallocate(b)
+ deallocate(c)
+end program map_multi
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
index 445cbabb8ca4a0ddfdbebb6638ea86341561d631..1d97dd382d49b2a8a2c2f00260739f40a50f3fa0 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
@@ -21,15 +21,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
index 7b206ac20428e99e989095474d74920b36c94f11..4307f50c46e97119826f7e441175e53bf5e15851 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
@@ -1,9 +1,6 @@
! { dg-do run }
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-/* Nullify the 'finalize' clause, which disturbs reference counting. */
+/* Nullify the 'finalize' clause. */
#define finalize
#include "mdc-refcount-1-1-1.f90"
-
-! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
index 8554534b2f271cc48565f9beeaf48238bcf1fb38..e6f3f4afc3b89f340d53d054b9f9b60fbbb43e57 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
@@ -23,15 +23,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
index 8e696cc70e80881ca94b6edf77b4ce2ee6103a9d..78f54e64dcea6e62bcb0563c8d737d773bbf81ce 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
@@ -23,15 +23,7 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a)
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
index 070a6f8e1493fbe8772a92c6d858b61502715f21..f9dcb485b2605e3154cd4fe8601bc7fcc29ef827 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
@@ -24,15 +24,7 @@ program main
if (.not. acc_is_present(var)) stop 2
!$acc exit data detach(var%a)
- print *, "CheCKpOInT1"
- ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
!$acc exit data delete(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
- !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
- !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
- print *, "CheCKpOInT2"
- ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
index b22e411567ffb9a649624e32afb92001a52bf8ae..fbd52373946baffab1ce38536a57e9e6cb364a95 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
@@ -23,16 +23,15 @@ program main
if (.not. acc_is_present(var%a)) stop 1
if (.not. acc_is_present(var)) stop 2
+ !$acc exit data detach(var%a) finalize
print *, "CheCKpOInT1"
! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
- !$acc exit data detach(var%a) finalize
- !TODO goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
- !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+ !$acc exit data delete(var%a)
+ !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
!TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
!TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
print *, "CheCKpOInT2"
! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
- !$acc exit data delete(var%a)
if (acc_is_present(var%a)) stop 3
if (.not. acc_is_present(var)) stop 4