libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c

   1 /* This test is similar to data-2.c, but it uses acc_* library functions
   2    to move data.  */
   3
   4 /* { dg-do run } */
   5
   6 #include <stdlib.h>
   7 #include <assert.h>
   8 #include <openacc.h>
   9
  10 int
  11 main (int argc, char **argv)
  12 {
  13   int N = 128; //1024 * 1024;
  14   float *a, *b, *c, *d, *e;
  15   void *d_a, *d_b, *d_c, *d_d;
  16   int i;
  17   int nbytes;
  18
  19   nbytes = N * sizeof (float);
  20
  21   a = (float *) malloc (nbytes);
  22   b = (float *) malloc (nbytes);
  23   c = (float *) malloc (nbytes);
  24   d = (float *) malloc (nbytes);
  25   e = (float *) malloc (nbytes);
  26
  27   for (i = 0; i < N; i++)
  28     {
  29       a[i] = 3.0;
  30       b[i] = 0.0;
  31     }
  32
  33   d_a = acc_copyin (a, nbytes);
  34   d_b = acc_copyin (b, nbytes);
  35   acc_copyin (&N, sizeof (int));
  36
  37 #pragma acc parallel present (a[0:N], b[0:N], N) async wait
  38 #pragma acc loop
  39   for (i = 0; i < N; i++)
  40     b[i] = a[i];
  41
  42   acc_wait_all ();
  43
  44   acc_memcpy_from_device (a, d_a, nbytes);
  45   acc_memcpy_from_device (b, d_b, nbytes);
  46
  47   for (i = 0; i < N; i++)
  48     {
  49       assert (a[i] == 3.0);
  50       assert (b[i] == 3.0);
  51     }
  52
  53   for (i = 0; i < N; i++)
  54     {
  55       a[i] = 2.0;
  56       b[i] = 0.0;
  57     }
  58
  59   acc_update_device (a, nbytes);
  60   acc_update_device (b, nbytes);
  61
  62 #pragma acc parallel present (a[0:N], b[0:N], N)  async (1)
  63 #pragma acc loop
  64   for (i = 0; i < N; i++)
  65     b[i] = a[i];
  66
  67   acc_wait (1);
  68
  69   acc_memcpy_from_device (a, d_a, nbytes);
  70   acc_memcpy_from_device (b, d_b, nbytes);
  71
  72   for (i = 0; i < N; i++)
  73     {
  74       assert (a[i] == 2.0);
  75       assert (b[i] == 2.0);
  76     }
  77
  78   for (i = 0; i < N; i++)
  79     {
  80       a[i] = 3.0;
  81       b[i] = 0.0;
  82       c[i] = 0.0;
  83       d[i] = 0.0;
  84     }
  85
  86   acc_update_device (a, nbytes);
  87   acc_update_device (b, nbytes);
  88   d_c = acc_copyin (c, nbytes);
  89   d_d = acc_copyin (d, nbytes);
  90
  91 #pragma acc parallel present (a[0:N], b[0:N], N) async (1)
  92 #pragma acc loop
  93   for (i = 0; i < N; i++)
  94     b[i] = (a[i] * a[i] * a[i]) / a[i];
  95
  96 #pragma acc parallel present (a[0:N], c[0:N], N) async (2)
  97 #pragma acc loop
  98   for (i = 0; i < N; i++)
  99     c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
 100
 101 #pragma acc parallel present (a[0:N], d[0:N], N) async (3)
 102 #pragma acc loop
 103   for (i = 0; i < N; i++)
 104     d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 105
 106   acc_wait_all ();
 107
 108   acc_memcpy_from_device (a, d_a, nbytes);
 109   acc_memcpy_from_device (b, d_b, nbytes);
 110   acc_memcpy_from_device (c, d_c, nbytes);
 111   acc_memcpy_from_device (d, d_d, nbytes);
 112
 113   for (i = 0; i < N; i++)
 114     {
 115       if (a[i] != 3.0)
 116         abort ();
 117
 118       if (b[i] != 9.0)
 119         abort ();
 120
 121       if (c[i] != 4.0)
 122         abort ();
 123
 124       if (d[i] != 1.0)
 125         abort ();
 126     }
 127
 128   for (i = 0; i < N; i++)
 129     {
 130       a[i] = 2.0;
 131       b[i] = 0.0;
 132       c[i] = 0.0;
 133       d[i] = 0.0;
 134       e[i] = 0.0;
 135     }
 136
 137   acc_update_device (a, nbytes);
 138   acc_update_device (b, nbytes);
 139   acc_update_device (c, nbytes);
 140   acc_update_device (d, nbytes);
 141   acc_copyin (e, nbytes);
 142
 143 #pragma acc parallel present (a[0:N], b[0:N], N) async (1)
 144   for (int ii = 0; ii < N; ii++)
 145     b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 146
 147 #pragma acc parallel present (a[0:N], c[0:N], N) async (2)
 148   for (int ii = 0; ii < N; ii++)
 149     c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 150
 151 #pragma acc parallel present (a[0:N], d[0:N], N) async (3)
 152   for (int ii = 0; ii < N; ii++)
 153     d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 154
 155 #pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \
 156   async (4)
 157   for (int ii = 0; ii < N; ii++)
 158     e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 159
 160   acc_wait_all ();
 161   acc_copyout (a, nbytes);
 162   acc_copyout (b, nbytes);
 163   acc_copyout (c, nbytes);
 164   acc_copyout (d, nbytes);
 165   acc_copyout (e, nbytes);
 166   acc_delete (&N, sizeof (int));
 167
 168   for (i = 0; i < N; i++)
 169     {
 170       if (a[i] != 2.0)
 171         abort ();
 172
 173       if (b[i] != 4.0)
 174         abort ();
 175
 176       if (c[i] != 4.0)
 177         abort ();
 178
 179       if (d[i] != 1.0)
 180         abort ();
 181
 182       if (e[i] != 11.0)
 183         abort ();
 184     }
 185
 186   return 0;
 187 }