diff options
Diffstat (limited to 'libgomp')
31 files changed, 449 insertions, 52 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 6a39087..7ad7ff4 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,48 @@ +2016-05-24 Cesar Philippidis <cesar@codesourcery.com> + + * oacc-mem.c (acc_malloc): Update handling of shared-memory targets. + (acc_free): Likewise. + (acc_memcpy_to_device): Likewise. + (acc_memcpy_from_device): Likewise. + (acc_deviceptr): Likewise. + (acc_hostptr): Likewise. + (acc_is_present): Likewise. + (acc_map_data): Likewise. + (acc_unmap_data): Likewise. + (present_create_copy): Likewise. + (delete_copyout): Likewise. + (update_dev_host): Likewise. + * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Remove xfail. + * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: New test. + * testsuite/libgomp.oacc-c-c++-common/data-2.c: Adjust test. + * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: New test. + * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Adjust test so that + it only runs on nvptx targets. + * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise. + 2016-05-23 Martin Jambor <mjambor@suse.cz> * testsuite/libgomp.hsa.c/switch-sbr-2.c: New test. diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index ce1905c..665e208 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -32,6 +32,7 @@ #include "gomp-constants.h" #include "oacc-int.h" #include <stdint.h> +#include <string.h> #include <assert.h> /* Return block containing [H->S), or NULL if not contained. The device lock @@ -104,6 +105,9 @@ acc_malloc (size_t s) assert (thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return malloc (s); + return thr->dev->alloc_func (thr->dev->target_id, s); } @@ -124,6 +128,9 @@ acc_free (void *d) struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return free (d); + gomp_mutex_lock (&acc_dev->lock); /* We don't have to call lazy open here, as the ptr value must have @@ -154,6 +161,12 @@ acc_memcpy_to_device (void *d, void *h, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (d, h, s); + return; + } + thr->dev->host2dev_func (thr->dev->target_id, d, h, s); } @@ -166,6 +179,12 @@ acc_memcpy_from_device (void *h, void *d, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (h, d, s); + return; + } + thr->dev->dev2host_func (thr->dev->target_id, h, d, s); } @@ -184,6 +203,9 @@ acc_deviceptr (void *h) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&dev->lock); n = lookup_host (dev, h, 1); @@ -218,6 +240,9 @@ acc_hostptr (void *d) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return d; + gomp_mutex_lock (&acc_dev->lock); n = lookup_dev (acc_dev->openacc.data_environ, d, 1); @@ -252,6 +277,9 @@ acc_is_present (void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h != NULL; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -271,7 +299,7 @@ acc_is_present (void *h, size_t s) void acc_map_data (void *h, void *d, size_t s) { - struct target_mem_desc *tgt; + struct target_mem_desc *tgt = NULL; size_t mapnum = 1; void *hostaddrs = h; void *devaddrs = d; @@ -287,9 +315,6 @@ acc_map_data (void *h, void *d, size_t s) { if (d != h) gomp_fatal ("cannot map data on shared-memory system"); - - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); } else { @@ -335,6 +360,10 @@ acc_unmap_data (void *h) /* No need to call lazy open, as the address must have been mapped. */ + /* This is a no-op on shared-memory targets. */ + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + size_t host_size; gomp_mutex_lock (&acc_dev->lock); @@ -405,6 +434,9 @@ present_create_copy (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -496,6 +528,9 @@ delete_copyout (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -553,6 +588,9 @@ update_dev_host (int is_dev, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c index f3b490a..d478ce2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@ -1,6 +1,4 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* <http://news.gmane.org/find-root.php?message_id=%3C87pp0aaksc.fsf%40kepler.schwinge.homeip.net%3E>. - { dg-xfail-run-if "TODO" { *-*-* } } */ /* { dg-additional-options "-lcuda" } */ #include <openacc.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c new file mode 100644 index 0000000..e1aa2c9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c @@ -0,0 +1,185 @@ +/* This test is similar to data-2.c, but it uses acc_* library functions + to move data. */ + +/* { dg-do run } */ + +#include <stdlib.h> +#include <assert.h> +#include <openacc.h> + +int +main (int argc, char **argv) +{ + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + void *d_a, *d_b, *d_c, *d_d; + int i; + int nbytes; + + nbytes = N * sizeof (float); + + a = (float *) malloc (nbytes); + b = (float *) malloc (nbytes); + c = (float *) malloc (nbytes); + d = (float *) malloc (nbytes); + e = (float *) malloc (nbytes); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + d_a = acc_copyin (a, nbytes); + d_b = acc_copyin (b, nbytes); + acc_copyin (&N, sizeof (int)); + +#pragma acc parallel present (a[0:N], b[0:N], N) async wait +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 3.0); + assert (b[i] == 3.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 2.0); + assert (b[i] == 2.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + d_c = acc_copyin (c, nbytes); + d_d = acc_copyin (d, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = (a[i] * a[i] * a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) +#pragma acc loop + for (i = 0; i < N; i++) + c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) +#pragma acc loop + for (i = 0; i < N; i++) + d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + acc_memcpy_from_device (c, d_c, nbytes); + acc_memcpy_from_device (d, d_d, nbytes); + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 9.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + acc_update_device (c, nbytes); + acc_update_device (d, nbytes); + acc_copyin (e, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) + for (int ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) + for (int ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) + for (int ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \ + async (4) + for (int ii = 0; ii < N; ii++) + e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; + + acc_wait_all (); + acc_copyout (a, nbytes); + acc_copyout (b, nbytes); + acc_copyout (c, nbytes); + acc_copyout (d, nbytes); + acc_copyout (e, nbytes); + acc_delete (&N, sizeof (int)); + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 4.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + + if (e[i] != 11.0) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c index f867a66..c1c0825 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions. */ + /* { dg-do run } */ #include <stdlib.h> @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -76,17 +78,17 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -120,26 +122,27 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index 747109f..0bf706a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions with 'acc update'. */ + /* { dg-do run } */ #include <stdlib.h> @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc update device (a[0:N], b[0:N]) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -78,17 +80,17 @@ main (int argc, char **argv) #pragma acc update device (b[0:N]) async (2) #pragma acc enter data copyin (c[0:N], d[0:N]) async (3) -#pragma acc parallel async (1) wait (1,2) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1,2) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1,3) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1,3) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -123,27 +125,28 @@ main (int argc, char **argv) #pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) #pragma acc enter data copyin (e[0:N]) async (5) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1,5) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1,5) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc exit data delete (N) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) @@ -162,5 +165,11 @@ main (int argc, char **argv) abort (); } + free (a); + free (b); + free (c); + free (d); + free (e); + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c new file mode 100644 index 0000000..b5b37b2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c @@ -0,0 +1,70 @@ +/* Verify enter/exit data interoperablilty between pragmas and + acc library calls. */ + +/* { dg-do run } */ + +#include <stdlib.h> +#include <assert.h> +#include <openacc.h> + +int +main () +{ + int *p = (int *)malloc (sizeof (int)); + + /* Test 1: pragma input, library output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 1; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 1); + + /* Test 2: library input, pragma output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 2; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 2); + + /* Test 3: library input, library output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 3); + + /* Test 4: pragma input, pragma output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 3); + + free (p); + + return 0; +} 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 7098ef3..d665533 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdlib.h> #include <openacc.h> 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 a9632f7..ee21257 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdlib.h> #include <openacc.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c index 4f6a731..50c1701 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_copyout. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdlib.h> #include <openacc.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c index 28e4e5c..c81a78d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Test if duplicate data mappings with acc_copy_in. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c index 7d1767e..a3487e8 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c @@ -1,4 +1,7 @@ -/* { dg-do run } */ +/* Check acc_copyout failure with acc_device_nvidia. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + #include <stdio.h> #include <stdlib.h> 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 160b33c..b686cc9 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Verify that acc_delete unregisters data mappings on the device. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c index 4f8e14c..25ceb3a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c index d908700..b170f81 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c index a6c0197..65ff440 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c index 2339dd6..fd4dc59 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c index d7de8e3..09e2817 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create, acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdlib.h> #include <openacc.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c index bb709d3..5f00ccb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create and acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c index 9304daa..7a96ab26 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with a NULL address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c index 92e3858..318a060 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with size zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c index e81627d..9bc9ecc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid partial acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c index 031c731..a24916d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid acc_present_or_create on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c index de5d1c1..30b90d4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device on unmapped data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c index 0d593f0..5db2912 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with a NULL data address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c index e98ecc4..8bbf016 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with size zero data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c index f26fc33..c214042 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <string.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c index 253ce59..afa137f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a size zero data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <string.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c index cfbb077..25c70c2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c index 5de376d..a8ee7df 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c index 3e621c3..fc221f4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with data size of zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include <stdio.h> #include <stdlib.h> |