[OpenACC 7/7] host_data construct (runtime tests)

[James Norris <jnorris@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 295 bytes --]


     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.


[-- Attachment #2: p7.patch --]
[-- Type: text/x-patch, Size: 4123 bytes --]

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;
+}