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

[PR92840] [OpenACC] Refuse 'acc_unmap_data' unless mapped by 'acc_map_data'

Browse source 
libgomp/
	PR libgomp/92840
	* oacc-mem.c (acc_map_data): Clarify reference counting behavior.
	(acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Adjust.

From-SVN: r279145
This commit is contained in:
Thomas Schwinge 2019-12-09 23:52:36 +01:00 • committed by Thomas Schwinge
parent 49070d0670
commit e103542bc8
7 changed files with 126 additions and 17 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

12
libgomp/ChangeLog
View file

@ -1,5 +1,17 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
PR libgomp/92840
* oacc-mem.c (acc_map_data): Clarify reference counting behavior.
(acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'.
* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c:
New file.
* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Adjust.
PR libgomp/92511
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: Remove
this file...

18
libgomp/oacc-mem.c
View file

@ -407,7 +407,11 @@ acc_map_data (void *h, void *d, size_t s)
tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
&kinds, true, GOMP_MAP_VARS_OPENACC);
tgt->list[0].key->refcount = REFCOUNT_INFINITY;
splay_tree_key n = tgt->list[0].key;
assert (n->refcount == 1);
assert (n->dynamic_refcount == 0);
/* Special reference counting behavior. */
n->refcount = REFCOUNT_INFINITY;
if (profiling_p)
{
@ -459,6 +463,18 @@ acc_unmap_data (void *h)
gomp_fatal ("[%p,%d] surrounds %p",
(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
the different 'REFCOUNT_INFINITY' cases, or simply separate
'REFCOUNT_INFINITY' values per different usage ('REFCOUNT_ACC_MAP_DATA'
etc.)? */
else if (n->refcount != REFCOUNT_INFINITY)
{
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("refusing to unmap block [%p,+%d] that has not been mapped"
" by 'acc_map_data'",
(void *) h, (int) host_size);
}
/* Mark for removal. */
n->refcount = 1;

27
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c Normal file
View file

@ -0,0 +1,27 @@
/* Verify that we refuse 'acc_unmap_data', after 'acc_create'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 101;
char *h = (char *) malloc (N);
void *d = acc_create (h, N - 3);
if (!d)
abort ();
fprintf (stderr, "CheCKpOInT\n");
acc_unmap_data (h);
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+98\\\] that has not been mapped by 'acc_map_data'" } */
/* { dg-shouldfail "" } */

25
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c Normal file
View file

@ -0,0 +1,25 @@
/* Verify that we refuse 'acc_unmap_data', after '#pragma acc enter data create'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 101;
char *h = (char *) malloc (N);
#pragma acc enter data create (h[0:N - 77])
fprintf (stderr, "CheCKpOInT\n");
acc_unmap_data (h);
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+24\\\] that has not been mapped by 'acc_map_data'" } */
/* { dg-shouldfail "" } */

26
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c Normal file
View file

@ -0,0 +1,26 @@
/* Verify that we refuse 'acc_unmap_data', inside 'data'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 101;
char *h = (char *) malloc (N);
#pragma acc data create (h[0:N - 55])
{
fprintf (stderr, "CheCKpOInT\n");
acc_unmap_data (h);
}
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
/* { dg-output "refusing to unmap block \\\[\[0-9a-fA-FxX\]+,\\\+46\\\] that has not been mapped by 'acc_map_data'" } */
/* { dg-shouldfail "" } */

21
libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
View file

@ -266,13 +266,15 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
d = (float *) acc_deviceptr (&a[0]);
acc_unmap_data (&a[0]);
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
d = (float *) acc_deviceptr (&b[0]);
acc_unmap_data (&b[0]);
acc_free (d);
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)
{
@ -475,11 +477,10 @@ main (int argc, char **argv)
abort ();
}
d = (float *) acc_deviceptr (a);
acc_delete (a, N * sizeof (float));
acc_unmap_data (a);
acc_free (d);
if (acc_is_present (a, N * sizeof (float)))
abort ();
d = (float *) acc_deviceptr (c);

14
libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
View file

@ -290,13 +290,15 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
d = (float *) acc_deviceptr (&a[0]);
acc_unmap_data (&a[0]);
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
d = (float *) acc_deviceptr (&b[0]);
acc_unmap_data (&b[0]);
acc_free (d);
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)