This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[OpenACC 7/7] host_data construct (runtime tests)
- From: James Norris <jnorris at codesourcery dot com>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc: "Joseph S. Myers" <joseph at codesourcery dot com>, Nathan Sidwell <Nathan_Sidwell at mentor dot com>, Jakub Jelinek <jakub at redhat dot com>
- Date: Thu, 22 Oct 2015 14:20:24 -0500
- Subject: [OpenACC 7/7] host_data construct (runtime tests)
- Authentication-results: sourceware.org; auth=none
- References: <56293476 dot 5020801 at codesourcery dot com>
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
- New runtime tests for host_data.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
new file mode 100644
index 0000000..15ccb27
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
@@ -0,0 +1,125 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-lcuda -lcublas -lcudart" } */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+#include <cuda.h>
+#include <cuda_runtime_api.h>
+#include <cublas_v2.h>
+
+void
+saxpy_host (int n, float a, float *x, float *y)
+{
+ int i;
+
+ for (i = 0; i < n; i++)
+ y[i] = y[i] + a * x[i];
+}
+
+#pragma acc routine
+void
+saxpy_target (int n, float a, float *x, float *y)
+{
+ int i;
+
+ for (i = 0; i < n; i++)
+ y[i] = y[i] + a * x[i];
+}
+
+int
+main(int argc, char **argv)
+{
+ const int N = 8;
+ int i;
+ float *x_ref, *y_ref;
+ float *x, *y;
+ cublasHandle_t h;
+ float a = 2.0;
+
+ x_ref = (float*) malloc (N * sizeof(float));
+ y_ref = (float*) malloc (N * sizeof(float));
+
+ x = (float*) malloc (N * sizeof(float));
+ y = (float*) malloc (N * sizeof(float));
+
+#pragma acc data copyin (x[0:N]) copy (y[0:N])
+ {
+ float *xp, *yp;
+#pragma acc host_data use_device (x, y)
+ {
+#pragma acc parallel pcopy (xp, yp) present (x, y)
+ {
+ xp = x;
+ yp = y;
+ }
+ }
+
+ if (xp != acc_deviceptr (x) || yp != acc_deviceptr (y))
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ x[i] = x_ref[i] = 4.0 + i;
+ y[i] = y_ref[i] = 3.0;
+ }
+
+ saxpy_host (N, a, x_ref, y_ref);
+
+ cublasCreate (&h);
+
+#pragma acc data copyin (x[0:N]) copy (y[0:N])
+ {
+#pragma acc host_data use_device (x, y)
+ {
+ cublasSaxpy (h, N, &a, x, 1, y, 1);
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (y[i] != y_ref[i])
+ abort ();
+ }
+
+#pragma acc data create (x[0:N]) copyout (y[0:N])
+ {
+#pragma acc kernels
+ for (i = 0; i < N; i++)
+ y[i] = 3.0;
+
+#pragma acc host_data use_device (x, y)
+ {
+ cublasSaxpy (h, N, &a, x, 1, y, 1);
+ }
+ }
+
+ cublasDestroy (h);
+
+ for (i = 0; i < N; i++)
+ {
+ if (y[i] != y_ref[i])
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ y[i] = 3.0;
+
+#pragma acc data copyin (x[0:N]) copyin (a, N) copy (y[0:N])
+ {
+#pragma acc host_data use_device (x, y)
+ {
+#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a, N)
+ saxpy_target (N, a, x, y);
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (y[i] != y_ref[i])
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
new file mode 100644
index 0000000..511ec64
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
@@ -0,0 +1,50 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+struct by_lightning {
+ int a;
+ int b;
+ int c;
+};
+
+int main (int argc, char* argv[])
+{
+ int x;
+ void *q = NULL, *r = NULL, *p = NULL, *s = NULL, *t = NULL;
+ long u;
+ struct by_lightning on_the_head = {1, 2, 3};
+ int arr[10], *f = NULL;
+ _Complex float cf;
+ #pragma acc enter data copyin (x, arr, on_the_head, cf)
+ #pragma acc host_data use_device (x, arr, on_the_head, cf)
+ {
+ q = &x;
+ {
+ f = &arr[5];
+ r = f;
+ s = &__real__ cf;
+ t = &on_the_head.c;
+ u = (long) &__imag__ cf;
+ #pragma acc parallel copyout(p) present (x, arr, on_the_head, cf)
+ {
+ /* This will not (and must not) call GOACC_deviceptr, but '&x' will be
+ the address on the device (if appropriate) regardless. */
+ p = &x;
+ }
+ }
+ }
+ #pragma acc exit data delete (x)
+
+#if ACC_MEM_SHARED
+ if (q != &x || f != &arr[5] || r != f || s != &(__real__ cf)
+ || t != &on_the_head.c || u != (long) &(__imag__ cf) || p != &x)
+ abort ();
+#else
+ if (q == &x || f == &arr[5] || r != f || s == &(__real__ cf)
+ || t == &on_the_head.c || u == (long) &(__imag__ cf) || p == &x)
+ abort ();
+#endif
+
+ return 0;
+}