Hi! On 2018-05-25T13:01:58-0700, Cesar Philippidis <ce...@codesourcery.com> wrote: > This patch updates GCC's to support OpenACC 2.5's data clause semantics.
Per <https://gcc.gnu.org/PR92970> "OpenACC 2.5: 'acc_delete' etc. on non-present data is a no-op", which this patch didn't address. I wanted to delay fixing this until I got the intended OpenACC 2.6 ff. semantics clarified with the OpenACC Technical Committee, but it turned out that fixing this now would be useful for other reasons, so see attached "[PR92726, PR92970, PR92984] [OpenACC] Clarify 'acc_delete' etc. for 'NULL'-in, non-present data, or size zero"; committed to trunk in r279532. More C/C++ and also Fortran test cases (that exercises all the different code paths that we have in 'libgomp/oacc-mem.c:GOACC_enter_exit_data', related to 'find_pointer' handling etc.) shall then follow later (no hurry with that). Grüße Thomas
From f7b1686558c2515511917aaeb74269b7e85ae09b Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 18 Dec 2019 17:01:11 +0000
Subject: [PATCH] [PR92726, PR92970, PR92984] [OpenACC] Clarify 'acc_delete'
etc. for 'NULL'-in, non-present data, or size zero
PR92970 "OpenACC 2.5: 'acc_delete' etc. on non-present data is a no-op" is an
actual bug fix, and the other ones are fall-out, currently undefined behavior.
libgomp/
PR libgomp/92726
PR libgomp/92970
PR libgomp/92984
* oacc-mem.c (delete_copyout): No-op behavior if 'lookup_host'
fails.
(GOACC_enter_exit_data): Simplify accordingly.
* testsuite/libgomp.oacc-c-c++-common/pr92970-1.c: New file,
subsuming...
* testsuite/libgomp.oacc-c-c++-common/lib-17.c: ... this file...
* testsuite/libgomp.oacc-c-c++-common/lib-18.c: ..., and this
file.
* testsuite/libgomp.oacc-c-c++-common/pr92984-1.c: New file,
subsuming...
* testsuite/libgomp.oacc-c-c++-common/lib-21.c: ... this file...
* testsuite/libgomp.oacc-c-c++-common/lib-29.c: ..., and this
file.
* testsuite/libgomp.oacc-c-c++-common/pr92726-1.c: New file,
subsuming...
* testsuite/libgomp.oacc-c-c++-common/lib-28.c: ... this file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279532 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 20 ++++
libgomp/oacc-mem.c | 28 ++---
.../libgomp.oacc-c-c++-common/lib-17.c | 38 -------
.../libgomp.oacc-c-c++-common/lib-18.c | 38 -------
.../libgomp.oacc-c-c++-common/lib-21.c | 35 ------
.../libgomp.oacc-c-c++-common/lib-28.c | 32 ------
.../libgomp.oacc-c-c++-common/lib-29.c | 32 ------
.../libgomp.oacc-c-c++-common/pr92726-1.c | 26 +++++
.../libgomp.oacc-c-c++-common/pr92970-1.c | 33 ++++++
.../libgomp.oacc-c-c++-common/pr92984-1.c | 100 ++++++++++++++++++
10 files changed, 190 insertions(+), 192 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92726-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92970-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92984-1.c
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index c4283fdfe1d..871a1537c77 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,25 @@
2019-12-18 Thomas Schwinge <tho...@codesourcery.com>
+ PR libgomp/92726
+ PR libgomp/92970
+ PR libgomp/92984
+ * oacc-mem.c (delete_copyout): No-op behavior if 'lookup_host'
+ fails.
+ (GOACC_enter_exit_data): Simplify accordingly.
+ * testsuite/libgomp.oacc-c-c++-common/pr92970-1.c: New file,
+ subsuming...
+ * testsuite/libgomp.oacc-c-c++-common/lib-17.c: ... this file...
+ * testsuite/libgomp.oacc-c-c++-common/lib-18.c: ..., and this
+ file.
+ * testsuite/libgomp.oacc-c-c++-common/pr92984-1.c: New file,
+ subsuming...
+ * testsuite/libgomp.oacc-c-c++-common/lib-21.c: ... this file...
+ * testsuite/libgomp.oacc-c-c++-common/lib-29.c: ..., and this
+ file.
+ * testsuite/libgomp.oacc-c-c++-common/pr92726-1.c: New file,
+ subsuming...
+ * testsuite/libgomp.oacc-c-c++-common/lib-28.c: ... this file.
+
* oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data'
'finalize' handling.
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b21d83c37d8..32bf3656029 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -659,7 +659,9 @@ acc_pcopyin (void *h, size_t s)
static void
delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
{
- splay_tree_key n;
+ /* No need to call lazy open, as the data must already have been
+ mapped. */
+
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
@@ -677,16 +679,10 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
gomp_mutex_lock (&acc_dev->lock);
- n = lookup_host (acc_dev, h, s);
-
- /* No need to call lazy open, as the data must already have been
- mapped. */
-
+ splay_tree_key n = lookup_host (acc_dev, h, s);
if (!n)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("[%p,%d] is not mapped", (void *)h, (int)s);
- }
+ /* PR92726, RP92970, PR92984: no-op. */
+ goto out;
if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end)
{
@@ -741,6 +737,7 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
}
}
+ out:
gomp_mutex_unlock (&acc_dev->lock);
if (profiling_p)
@@ -1224,13 +1221,10 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
{
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
- if (acc_is_present (hostaddrs[i], sizes[i]))
- {
- if (finalize)
- acc_delete_finalize_async (hostaddrs[i], sizes[i], async);
- else
- acc_delete_async (hostaddrs[i], sizes[i], async);
- }
+ if (finalize)
+ acc_delete_finalize_async (hostaddrs[i], sizes[i], async);
+ else
+ acc_delete_async (hostaddrs[i], sizes[i], async);
break;
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
deleted file mode 100644
index a3487e8f5bf..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
+++ /dev/null
@@ -1,38 +0,0 @@
-/* Check acc_copyout failure with acc_device_nvidia. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- (void) acc_copyin (h, N);
-
- acc_copyout (h, N);
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyout (h, N);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
deleted file mode 100644
index 93bfb99f415..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
+++ /dev/null
@@ -1,38 +0,0 @@
-/* Verify that acc_delete unregisters data mappings on the device. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d = acc_copyin (h, N);
-
- acc_delete (h, N);
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyout (h, N);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
deleted file mode 100644
index b170f81229c..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
+++ /dev/null
@@ -1,35 +0,0 @@
-/* Exercise acc_copyin and acc_copyout on nvidia targets. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- (void) acc_copyin (h, N);
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_copyout (h, 0);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,0\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
deleted file mode 100644
index 7a96ab26ebd..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
+++ /dev/null
@@ -1,32 +0,0 @@
-/* Exercise acc_delete with a NULL address on nvidia targets. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- unsigned char *h;
- void *d;
-
- h = (unsigned char *) malloc (N);
-
- d = acc_create (h, N);
- if (!d)
- abort ();
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_delete (0, N);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[^\n\r]*,256\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
deleted file mode 100644
index 318a060f228..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
+++ /dev/null
@@ -1,32 +0,0 @@
-/* Exercise acc_delete with size zero on nvidia targets. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- unsigned char *h;
- void *d;
-
- h = (unsigned char *) malloc (N);
-
- d = acc_create (h, N);
- if (!d)
- abort ();
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_delete (h, 0);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,0\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92726-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92726-1.c
new file mode 100644
index 00000000000..fb69adf4c40
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92726-1.c
@@ -0,0 +1,26 @@
+/* Verify that 'acc_delete' etc. with a 'NULL' address is a no-op. */
+
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main (int argc, char **argv)
+{
+ const int N = 256;
+
+ unsigned char *a = (unsigned char *) malloc (N);
+ assert (a);
+
+ void *a_d = acc_create (a, N);
+ assert (a_d);
+
+ acc_delete (NULL, N);
+ assert (acc_is_present (a, N));
+ //TODO similar for others.
+
+ acc_delete (a, N);
+ free (a);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92970-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92970-1.c
new file mode 100644
index 00000000000..380f6793454
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92970-1.c
@@ -0,0 +1,33 @@
+/* Verify that 'acc_delete' etc. on non-present data is a no-op. */
+
+#include <openacc.h>
+
+int
+main ()
+{
+ int a;
+
+ int async = 0;
+
+#pragma acc exit data copyout (a)
+ acc_copyout (&a, sizeof a);
+#pragma acc exit data copyout (a) async (async++)
+ acc_copyout_async (&a, sizeof a, async++);
+#pragma acc exit data copyout (a) finalize
+ acc_copyout_finalize (&a, sizeof a);
+#pragma acc exit data copyout (a) finalize async (async++)
+ acc_copyout_finalize_async (&a, sizeof a, async++);
+
+#pragma acc exit data delete (a)
+ acc_delete (&a, sizeof a);
+#pragma acc exit data delete (a) async (async++)
+ acc_delete_async (&a, sizeof a, async++);
+#pragma acc exit data delete (a) finalize
+ acc_delete_finalize (&a, sizeof a);
+#pragma acc exit data delete (a) finalize async (async++)
+ acc_delete_finalize_async (&a, sizeof a, async++);
+
+ acc_wait_all ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92984-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92984-1.c
new file mode 100644
index 00000000000..319d6ccfd35
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92984-1.c
@@ -0,0 +1,100 @@
+/* Verify that 'acc_delete' etc. with zero size is a no-op. */
+
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+
+#define UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+
+
+static void
+verify_mapped_unchanged (unsigned char *a, size_t N)
+{
+ assert (acc_is_present (a, N));
+
+ for (size_t i = 0; i < N; ++i)
+ assert (a[i] == (unsigned char) i);
+}
+
+int
+main (int argc, char **argv)
+{
+ const size_t N = 256;
+
+ unsigned char *a = (unsigned char *) malloc (N);
+ assert (a);
+
+ for (size_t i = 0; i < N; ++i)
+ a[i] = 51;
+
+ void *a_d = acc_copyin (a, N);
+ assert (a_d);
+
+ for (size_t i = 0; i < N; ++i)
+ a[i] = i;
+
+ int async = 0;
+
+ const size_t size = 0;
+
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data copyout (a[0:size])
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_copyout (a, size);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data copyout (a[0:size]) async (async++)
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_copyout_async (a, size, async++);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data copyout (a[0:size]) finalize
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_copyout_finalize (a, size);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data copyout (a[0:size]) finalize async (async++)
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_copyout_finalize_async (a, size, async++);
+ verify_mapped_unchanged (a, N);
+
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data delete (a[0:size])
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_delete (a, size);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data delete (a[0:size]) async (async++)
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_delete_async (a, size, async++);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data delete (a[0:size]) finalize
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_delete_finalize (a, size);
+ verify_mapped_unchanged (a, N);
+#ifndef UNHANDLED_GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+#pragma acc exit data delete (a[0:size]) finalize async (async++)
+ verify_mapped_unchanged (a, N);
+#endif
+ acc_delete_finalize_async (a, size, async++);
+ verify_mapped_unchanged (a, N);
+
+ acc_wait_all ();
+
+ acc_delete (a, N);
+#if !ACC_MEM_SHARED
+ assert (!acc_is_present (a, N));
+#endif
+ free (a);
+
+ return 0;
+}
--
2.17.1
signature.asc
Description: PGP signature
