libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c

   1 /* acc_present_or_create, acc_present_or_copyin, etc.  */
   2 /* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*".  */
   3
   4 #include <stdbool.h>
   5 #include <stdlib.h>
   6 #include <openacc.h>
   7
   8 int
   9 main (int argc, char **argv)
  10 {
  11   int *h, *d;
  12   const int N = 10000;
  13   const int S = N * sizeof *h;
  14   bool shared_mem;
  15
  16   h = (int *) malloc (S);
  17   if (!h)
  18     abort ();
  19   for (int i = 0; i < N; ++i)
  20     h[i] = i + 0;
  21
  22   shared_mem = acc_is_present (h, S);
  23
  24   d = (int *) acc_present_or_create (h, S);
  25   if (!d)
  26     abort ();
  27   if (shared_mem)
  28     if (h != d)
  29       abort ();
  30   if (!acc_is_present (h, S))
  31     abort ();
  32
  33 #pragma acc parallel loop deviceptr (d)
  34   for (int i = 0; i < N; ++i)
  35     {
  36       d[i] = i + 1;
  37     }
  38
  39   for (int i = 0; i < N; ++i)
  40     {
  41       if (h[i] != i + (shared_mem ? 1 : 0))
  42         abort ();
  43       h[i] = i + 2;
  44     }
  45
  46   {
  47     int *d_ = (int *) acc_present_or_create (h, S);
  48     if (d_ != d)
  49       abort ();
  50   }
  51
  52 #pragma acc parallel loop deviceptr (d)
  53   for (int i = 0; i < N; ++i)
  54     {
  55       if (d[i] != i + (shared_mem ? 2 : 1))
  56         abort ();
  57       d[i] = i + 3;
  58     }
  59
  60   for (int i = 0; i < N; ++i)
  61     {
  62       if (h[i] != i + (shared_mem ? 3 : 2))
  63         abort ();
  64       h[i] = i + 4;
  65     }
  66
  67   {
  68     int *d_ = (int *) acc_pcreate (h, S);
  69     if (d_ != d)
  70       abort ();
  71   }
  72
  73 #pragma acc parallel loop deviceptr (d)
  74   for (int i = 0; i < N; ++i)
  75     {
  76       if (d[i] != i + (shared_mem ? 4 : 3))
  77         abort ();
  78       d[i] = i + 5;
  79     }
  80
  81   for (int i = 0; i < N; ++i)
  82     {
  83       if (h[i] != i + (shared_mem ? 5 : 4))
  84         abort ();
  85       h[i] = i + 6;
  86     }
  87
  88   {
  89     int *d_ = (int *) acc_present_or_copyin (h, S);
  90     if (d_ != d)
  91       abort ();
  92   }
  93
  94 #pragma acc parallel loop deviceptr (d)
  95   for (int i = 0; i < N; ++i)
  96     {
  97       if (d[i] != i + (shared_mem ? 6 : 5))
  98         abort ();
  99       d[i] = i + 7;
 100     }
 101
 102   for (int i = 0; i < N; ++i)
 103     {
 104       if (h[i] != i + (shared_mem ? 7 : 6))
 105         abort ();
 106       h[i] = i + 8;
 107     }
 108
 109   {
 110     int *d_ = (int *) acc_pcopyin (h, S);
 111     if (d_ != d)
 112       abort ();
 113   }
 114
 115 #pragma acc parallel loop deviceptr (d)
 116   for (int i = 0; i < N; ++i)
 117     {
 118       if (d[i] != i + (shared_mem ? 8 : 7))
 119         abort ();
 120       d[i] = i + 9;
 121     }
 122
 123   for (int i = 0; i < N; ++i)
 124     {
 125       if (h[i] != i + (shared_mem ? 9 : 8))
 126         abort ();
 127       h[i] = i + 10;
 128     }
 129
 130   acc_copyout_finalize (h, S);
 131   d = NULL;
 132   if (!shared_mem)
 133     if (acc_is_present (h, S))
 134       abort ();
 135
 136   for (int i = 0; i < N; ++i)
 137     {
 138       if (h[i] != i + (shared_mem ? 10 : 9))
 139         abort ();
 140     }
 141
 142   d = (int *) acc_pcopyin (h, S);
 143   if (!d)
 144     abort ();
 145   if (shared_mem)
 146     if (h != d)
 147       abort ();
 148   if (!acc_is_present (h, S))
 149     abort ();
 150
 151 #pragma acc parallel loop deviceptr (d)
 152   for (int i = 0; i < N; ++i)
 153     {
 154       if (d[i] != i + (shared_mem ? 10 : 9))
 155         abort ();
 156       d[i] = i + 11;
 157     }
 158
 159   for (int i = 0; i < N; ++i)
 160     {
 161       if (h[i] != i + (shared_mem ? 11 : 9))
 162         abort ();
 163       h[i] = i + 12;
 164     }
 165
 166   {
 167     int *d_ = (int *) acc_pcopyin (h, S);
 168     if (d_ != d)
 169       abort ();
 170   }
 171
 172 #pragma acc parallel loop deviceptr (d)
 173   for (int i = 0; i < N; ++i)
 174     {
 175       if (d[i] != i + (shared_mem ? 12 : 11))
 176         abort ();
 177       d[i] = i + 13;
 178     }
 179
 180   for (int i = 0; i < N; ++i)
 181     {
 182       if (h[i] != i + (shared_mem ? 13 : 12))
 183         abort ();
 184       h[i] = i + 14;
 185     }
 186
 187   {
 188     int *d_ = (int *) acc_pcreate (h, S);
 189     if (d_ != d)
 190       abort ();
 191   }
 192
 193 #pragma acc parallel loop deviceptr (d)
 194   for (int i = 0; i < N; ++i)
 195     {
 196       if (d[i] != i + (shared_mem ? 14 : 13))
 197         abort ();
 198       d[i] = i + 15;
 199     }
 200
 201   for (int i = 0; i < N; ++i)
 202     {
 203       if (h[i] != i + (shared_mem ? 15 : 14))
 204         abort ();
 205       h[i] = i + 16;
 206     }
 207
 208   {
 209     int *d_ = (int *) acc_pcreate (h, S);
 210     if (d_ != d)
 211       abort ();
 212   }
 213
 214 #pragma acc parallel loop deviceptr (d)
 215   for (int i = 0; i < N; ++i)
 216     {
 217       if (d[i] != i + (shared_mem ? 16 : 15))
 218         abort ();
 219       d[i] = i + 17;
 220     }
 221
 222   for (int i = 0; i < N; ++i)
 223     {
 224       if (h[i] != i + (shared_mem ? 17 : 16))
 225         abort ();
 226       h[i] = i + 18;
 227     }
 228
 229   acc_update_self (h, S);
 230   if (!acc_is_present (h, S))
 231     abort ();
 232
 233   for (int i = 0; i < N; ++i)
 234     {
 235       if (h[i] != i + (shared_mem ? 18 : 17))
 236         abort ();
 237     }
 238
 239   acc_delete_finalize (h, S);
 240   d = NULL;
 241   if (!shared_mem)
 242     if (acc_is_present (h, S))
 243       abort();
 244
 245   free (h);
 246
 247   return 0;
 248 }