FreeChainXenon/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

openacc: Adjust dynamic reference count semantics

Browse source 
This patch adjusts how dynamic reference counts work so that they match
the semantics of the source program more closely, instead of representing
"excess" reference counts beyond those that represent pointers in the
internal libgomp splay-tree data structure. This allows some corner
cases to be handled more gracefully.

2020-07-10  Julian Brown  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	libgomp/
	* libgomp.h (struct splay_tree_key_s): Change virtual_refcount to
	dynamic_refcount.
	(struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA.
	* oacc-mem.c (acc_map_data): Substitute virtual_refcount for
	dynamic_refcount.
	(acc_unmap_data): Update comment.
	(goacc_map_var_existing, goacc_enter_datum): Adjust for
	dynamic_refcount semantics.
	(goacc_exit_datum_1, goacc_exit_datum): Re-add some error checking.
	Adjust for dynamic_refcount semantics.
	(goacc_enter_data_internal): Implement "present" case of dynamic
	memory-map handling here.  Update "non-present" case for
	dynamic_refcount semantics.
	(goacc_exit_data_internal): Use goacc_exit_datum_1.
	* target.c (gomp_map_vars_internal): Remove
	GOMP_MAP_VARS_OPENACC_ENTER_DATA handling.  Update for dynamic_refcount
	handling.
	(gomp_unmap_vars_internal): Remove virtual_refcount handling.
	(gomp_load_image_to_device): Substitute dynamic_refcount for
	virtual_refcount.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAILs.
	* testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs and
	trace output.
	* testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Remove
	trace output.
	* testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90: New
	test.
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
	Remove stale comment.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Adjust XFAIL.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
This commit is contained in:
Julian Brown 2020-05-22 04:06:10 -07:00
parent cb7effde6f
commit 6f5b4b64d2
17 changed files with 315 additions and 178 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
libgomp/libgomp.h
View file

@ -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

236
libgomp/oacc-mem.c
View file

@ -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;
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);
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];
}
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;

36
libgomp/target.c
View file

@ -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 (pragma_kind == GOMP_MAP_VARS_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;
}
@ -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;

9
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
View file

@ -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);

31
libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-1.c Normal file
View file

@ -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;
}

31
libgomp/testsuite/libgomp.oacc-c-c++-common/refcounting-2.c Normal file
View file

@ -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;
}

34
libgomp/testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c Normal file
View file

@ -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;
}

2
libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
View file

@ -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);

5
libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
View file

@ -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)" }

8
libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
View file

@ -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

49
libgomp/testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90 Normal file
View file

@ -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

8
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
View file

@ -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

5
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
View file

@ -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)" }

8
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
View file

@ -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

8
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
View file

@ -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

8
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
View file

@ -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

7
libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
View file

@ -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