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

[PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free'

Browse source 
libgomp/
	PR libgomp/92503
	* oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.

From-SVN: r279146
This commit is contained in:
Thomas Schwinge 2019-12-09 23:52:47 +01:00 • committed by Thomas Schwinge
parent e103542bc8
commit cec41816c1
18 changed files with 242 additions and 30 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

25
libgomp/ChangeLog
View file

@ -1,5 +1,30 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
PR libgomp/92503
* oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New
file.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
PR libgomp/92840
* oacc-mem.c (acc_map_data): Clarify reference counting behavior.
(acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'.

17
libgomp/oacc-mem.c
View file

@ -121,9 +121,6 @@ acc_malloc (size_t s)
return res;
}
/* OpenACC 2.0a (3.2.16) doesn't specify what to do in the event
the device address is mapped. We choose to check if it mapped,
and if it is, to unmap it. */
void
acc_free (void *d)
{
@ -152,13 +149,15 @@ acc_free (void *d)
(unless you got that null from acc_malloc). */
if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1)))
{
void *offset;
offset = d - k->tgt->tgt_start + k->tgt_offset;
void *offset = d - k->tgt->tgt_start + k->tgt_offset;
void *h = k->host_start + offset;
size_t h_size = k->host_end - k->host_start;
gomp_mutex_unlock (&acc_dev->lock);
acc_unmap_data ((void *)(k->host_start + offset));
/* PR92503 "[OpenACC] Behavior of 'acc_free' if the memory space is still
used in a mapping". */
gomp_fatal ("refusing to free device memory space at %p that is still"
" mapped at [%p,+%d]",
d, h, (int) h_size);
}
else
gomp_mutex_unlock (&acc_dev->lock);

28
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c Normal file
View file

@ -0,0 +1,28 @@
/* Verify that we refuse 'acc_free', after 'acc_map_data'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 108;
char *h = (char *) malloc (N);
void *d = acc_malloc (N - 10);
if (!d)
abort ();
acc_map_data (h, d, N - 19);
fprintf (stderr, "CheCKpOInT\n");
acc_free (d);
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+89\\\]" }
{ dg-shouldfail "" } */

27
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c Normal file
View file

@ -0,0 +1,27 @@
/* Verify that we refuse 'acc_free', after 'acc_create'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 108;
char *h = (char *) malloc (N);
void *d = acc_create (h, N - 1);
if (!d)
abort ();
fprintf (stderr, "CheCKpOInT\n");
acc_free (d);
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+107\\\]" }
{ dg-shouldfail "" } */

28
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c Normal file
View file

@ -0,0 +1,28 @@
/* Verify that we refuse 'acc_free', inside 'host_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 = 108;
char *h = (char *) malloc (N);
#pragma acc enter data create (h[0:N - 2])
#pragma acc host_data use_device (h)
{
fprintf (stderr, "CheCKpOInT\n");
acc_free (h);
}
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+106\\\]" }
{ dg-shouldfail "" } */

28
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c Normal file
View file

@ -0,0 +1,28 @@
/* Verify that we refuse 'acc_free', 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 = 108;
char *h = (char *) malloc (N);
#pragma acc enter data create (h[0:N - 3])
void *d = acc_deviceptr (h);
if (!d)
abort ();
fprintf (stderr, "CheCKpOInT\n");
acc_free (d);
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+105\\\]" }
{ dg-shouldfail "" } */

31
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c Normal file
View file

@ -0,0 +1,31 @@
/* Verify that we refuse 'acc_free', inside 'host_data', inside 'data'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 108;
char *h = (char *) malloc (N);
#pragma acc data create (h[0:N - 44])
{
#pragma acc host_data use_device (h)
{
fprintf (stderr, "CheCKpOInT\n");
acc_free (h);
}
}
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
TODO PR92877
{ dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+64\\\]" { xfail *-*-* } }
{ dg-shouldfail "" } */

32
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c Normal file
View file

@ -0,0 +1,32 @@
/* Verify that we refuse 'acc_free', inside 'data'. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
int
main ()
{
const int N = 108;
char *h = (char *) malloc (N);
#pragma acc data create (h[0:N - 21])
{
void *d = acc_deviceptr (h);
if (!d)
abort ();
fprintf (stderr, "CheCKpOInT\n");
acc_free (d);
}
return 0;
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
TODO PR92877
{ dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+87\\\]" { xfail *-*-* } }
{ dg-shouldfail "" } */

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

@ -103,7 +103,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)
{
@ -162,7 +165,7 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
@ -557,7 +560,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)
{

6
libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c
View file

@ -172,13 +172,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
acc_delete (&h_X[0], N * sizeof (float));
acc_delete (&h_Y1[0], N * sizeof (float));
free (h_X);
free (h_Y1);
free (h_Y2);
acc_free (d_X);
acc_free (d_Y);
context_check (pctx);
s = cublasDestroy (h);

6
libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
View file

@ -182,13 +182,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
acc_delete (&h_X[0], N * sizeof (float));
acc_delete (&h_Y1[0], N * sizeof (float));
free (h_X);
free (h_Y1);
free (h_Y2);
acc_free (d_X);
acc_free (d_Y);
context_check (pctx);
s = cublasDestroy (h);

6
libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c
View file

@ -163,13 +163,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
acc_delete (&h_X[0], N * sizeof (float));
acc_delete (&h_Y1[0], N * sizeof (float));
free (h_X);
free (h_Y1);
free (h_Y2);
acc_free (d_X);
acc_free (d_Y);
context_check (pctx);
s = cublasDestroy (h);

6
libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
View file

@ -176,13 +176,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
acc_delete (&h_X[0], N * sizeof (float));
acc_delete (&h_Y1[0], N * sizeof (float));
free (h_X);
free (h_Y1);
free (h_Y2);
acc_free (d_X);
acc_free (d_Y);
context_check (pctx);
s = cublasDestroy (h);

2
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
View file

@ -51,7 +51,7 @@ main (int argc, char **argv)
if (acc_is_present (h, 0) != 0)
abort ();
acc_free (d);
acc_delete (h, N);
if (acc_is_present (h, 1) != 0)
abort ();

2
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
View file

@ -48,7 +48,7 @@ main (int argc, char **argv)
abort ();
}
acc_free (d);
acc_delete (h, N);
for (i = 0; i < N; i++)
{

2
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
View file

@ -23,7 +23,7 @@ main (int argc, char **argv)
d = acc_copyin (h, N);
acc_free (d);
acc_delete (h, N);
fprintf (stderr, "CheCKpOInT\n");
acc_copyout (h, N);

2
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c
View file

@ -72,6 +72,8 @@ main (int argc, char **argv)
if (async > (sync * 1.5))
abort ();
acc_unmap_data (h);
acc_free (d);
free (h);

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

@ -112,7 +112,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)
{
@ -177,7 +180,7 @@ main (int argc, char **argv)
if (!acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&b[0], N * sizeof (float));
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
@ -609,7 +612,10 @@ main (int argc, char **argv)
if (acc_is_present (&b[0], (N * sizeof (float))))
abort ();
acc_free (d);
acc_delete (&a[0], N * sizeof (float));
if (acc_is_present (&a[0], N * sizeof (float)))
abort ();
for (i = 0; i < N; i++)
{