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.
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;
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)
{
}
}
+ out:
gomp_mutex_unlock (&acc_dev->lock);
if (profiling_p)
{
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:
+++ /dev/null
-/* 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 "" } */
+++ /dev/null
-/* 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 "" } */
+++ /dev/null
-/* 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 "" } */
+++ /dev/null
-/* 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 "" } */
+++ /dev/null
-/* 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 "" } */
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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;
+}
projects / gcc.git / commitdiff