aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorThomas Schwinge <thomas@codesourcery.com>2019-12-18 18:01:11 +0100
committerThomas Schwinge <tschwinge@gcc.gnu.org>2019-12-18 18:01:11 +0100
commitddb25eb9ca373b293da3e8f2c1520dbb72271367 (patch)
treebbf3d9e62726618daaa67ae2c9631c498594f8d7
parent32128577aed11aa21f3227edc2276da677e97636 (diff)
downloadgcc-ddb25eb9ca373b293da3e8f2c1520dbb72271367.zip
gcc-ddb25eb9ca373b293da3e8f2c1520dbb72271367.tar.gz
gcc-ddb25eb9ca373b293da3e8f2c1520dbb72271367.tar.bz2
[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. From-SVN: r279532
-rw-r--r--libgomp/ChangeLog20
-rw-r--r--libgomp/oacc-mem.c28
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c38
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c38
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c35
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c32
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c32
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/pr92726-1.c26
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/pr92970-1.c33
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/pr92984-1.c100
10 files changed, 190 insertions, 192 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index c4283fd..871a153 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,25 @@
2019-12-18 Thomas Schwinge <thomas@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 b21d83c..32bf365 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 a3487e8..0000000
--- 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 93bfb99..0000000
--- 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 b170f81..0000000
--- 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 7a96ab26..0000000
--- 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 318a060..0000000
--- 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 0000000..fb69adf
--- /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 0000000..380f679
--- /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 0000000..319d6cc
--- /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;
+}