From cec41816c18f2857f8362825222ef4de0a6e596e Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Mon, 9 Dec 2019 23:52:47 +0100 Subject: [PATCH] [PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free' 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 --- libgomp/ChangeLog | 25 +++++++++++++++ libgomp/oacc-mem.c | 17 +++++----- .../acc_free-pr92503-1.c | 28 ++++++++++++++++ .../acc_free-pr92503-2.c | 27 ++++++++++++++++ .../acc_free-pr92503-3-2.c | 28 ++++++++++++++++ .../acc_free-pr92503-3.c | 28 ++++++++++++++++ .../acc_free-pr92503-4-2.c | 31 ++++++++++++++++++ .../acc_free-pr92503-4.c | 32 +++++++++++++++++++ .../libgomp.oacc-c-c++-common/clauses-1.c | 12 +++++-- .../libgomp.oacc-c-c++-common/context-1.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-2.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-3.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-4.c | 6 ++-- .../libgomp.oacc-c-c++-common/lib-13.c | 2 +- .../libgomp.oacc-c-c++-common/lib-14.c | 2 +- .../libgomp.oacc-c-c++-common/lib-18.c | 2 +- .../libgomp.oacc-c-c++-common/lib-91.c | 2 ++ .../libgomp.oacc-c-c++-common/nested-1.c | 12 +++++-- 18 files changed, 242 insertions(+), 30 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 7606f17825d..62092a2d765 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,30 @@ 2019-12-09 Thomas Schwinge + 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'. diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 480b9fbb71b..81ebddf7580 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c new file mode 100644 index 00000000000..4fc6068ba98 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c @@ -0,0 +1,28 @@ +/* Verify that we refuse 'acc_free', after 'acc_map_data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c new file mode 100644 index 00000000000..3f6a8e57174 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c @@ -0,0 +1,27 @@ +/* Verify that we refuse 'acc_free', after 'acc_create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c new file mode 100644 index 00000000000..9f4504809eb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c @@ -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 +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c new file mode 100644 index 00000000000..162083051cf --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c @@ -0,0 +1,28 @@ +/* Verify that we refuse 'acc_free', after '#pragma acc enter data create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c new file mode 100644 index 00000000000..bbf44319687 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c @@ -0,0 +1,31 @@ +/* Verify that we refuse 'acc_free', inside 'host_data', inside 'data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c new file mode 100644 index 00000000000..6212f9eae47 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c @@ -0,0 +1,32 @@ +/* Verify that we refuse 'acc_free', inside 'data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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 "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c index d36a2f1c304..b0a96348c3a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c @@ -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++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c index dabc7063c68..2e3b62ebbd2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c index 6a52f746dcb..6bdcfe7d429 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c index ccd276cd98f..8f14560ea8b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c index 71365e8ed32..b403a5cf5cb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c index d6655335e21..aca4c252091 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c @@ -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 (); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c index ee21257c9a5..de6d38b060c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c @@ -48,7 +48,7 @@ main (int argc, char **argv) abort (); } - acc_free (d); + acc_delete (h, N); for (i = 0; i < N; i++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c index b686cc94815..93bfb99f415 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c index e00ef4f7206..36fff089b83 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c @@ -72,6 +72,8 @@ main (int argc, char **argv) if (async > (sync * 1.5)) abort (); + acc_unmap_data (h); + acc_free (d); free (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c index 7ebfb8a562b..4c599cda4b3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c @@ -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++) { -- 2.30.2