libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c

   1 /* { dg-do run } */
   2
   3 /* Based on '../libgomp.oacc-fortran/asyncwait-1.f90'.  */
   4
   5 /* { dg-additional-options "--param=openacc-kernels=decompose" } */
   6
   7 /* { dg-additional-options "-fopt-info-all-omp" }
   8    { dg-additional-options "-foffload=-fopt-info-all-omp" } */
   9
  10 /* { dg-additional-options "--param=openacc-privatization=noisy" }
  11    { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
  12    Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
  13    { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
  14
  15 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
  16    passed to 'incr' may be unset, and in that case, it will be set to [...]",
  17    so to maintain compatibility with earlier Tcl releases, we manually
  18    initialize counter variables:
  19    { dg-line l_dummy[variable c_compute 0 c_loop_i 0] }
  20    { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
  21    "WARNING: dg-line var l_dummy defined, but not used".  */
  22
  23 #include <stdlib.h>
  24
  25 #define N 64
  26
  27 int
  28 main (void)
  29 {
  30   int *a, *b, *c, *d, *e;
  31
  32   a = (int*)malloc (N * sizeof (*a));
  33   b = (int*)malloc (N * sizeof (*b));
  34   c = (int*)malloc (N * sizeof (*c));
  35   d = (int*)malloc (N * sizeof (*d));
  36   e = (int*)malloc (N * sizeof (*e));
  37
  38   for (int i = 0; i < N; ++i)
  39     {
  40       a[i] = 3;
  41       b[i] = 0;
  42     }
  43
  44 #pragma acc data copy (a[0:N]) copy (b[0:N])
  45   {
  46
  47 #pragma acc parallel async /* { dg-line l_compute[incr c_compute] } */
  48     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  49 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
  50     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
  51     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
  52     for (int i = 0; i < N; ++i)
  53       b[i] = a[i];
  54
  55 #pragma acc wait
  56   }
  57
  58   for (int i = 0; i < N; ++i)
  59     {
  60       if (a[i] != 3)
  61         abort ();
  62       if (b[i] != 3)
  63         abort ();
  64     }
  65
  66   for (int i = 0; i < N; ++i)
  67     {
  68       a[i] = 2;
  69       b[i] = 0;
  70     }
  71
  72 #pragma acc data copy (a[0:N]) copy (b[0:N])
  73   {
  74 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
  75     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  76 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
  77     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
  78     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
  79     for (int i = 0; i < N; ++i)
  80       b[i] = a[i];
  81
  82 #pragma acc wait (1)
  83   }
  84
  85   for (int i = 0; i < N; ++i)
  86     {
  87       if (a[i] != 2) abort ();
  88       if (b[i] != 2) abort ();
  89     }
  90
  91   for (int i = 0; i < N; ++i)
  92     {
  93       a[i] = 3;
  94       b[i] = 0;
  95       c[i] = 0;
  96       d[i] = 0;
  97     }
  98
  99 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
 100   {
 101
 102 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 103     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 104     for (int i = 0; i < N; ++i)
 105       b[i] = (a[i] * a[i] * a[i]) / a[i];
 106
 107 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 108     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 109     for (int i = 0; i < N; ++i)
 110       c[i] = (a[i] * 4) / a[i];
 111
 112
 113 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 114     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 115 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 116     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 117     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 118     for (int i = 0; i < N; ++i)
 119       d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 120
 121 #pragma acc wait (1)
 122   }
 123
 124   for (int i = 0; i < N; ++i)
 125     {
 126       if (a[i] != 3)
 127         abort ();
 128       if (b[i] != 9)
 129         abort ();
 130       if (c[i] != 4)
 131         abort ();
 132       if (d[i] != 1)
 133         abort ();
 134     }
 135
 136   for (int i = 0; i < N; ++i)
 137     {
 138       a[i] = 2;
 139       b[i] = 0;
 140       c[i] = 0;
 141       d[i] = 0;
 142       e[i] = 0;
 143     }
 144
 145 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
 146   {
 147
 148 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 149     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 150     for (int i = 0; i < N; ++i)
 151       b[i] = (a[i] * a[i] * a[i]) / a[i];
 152
 153 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 154     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 155 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 156     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 157     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 158     for (int i = 0; i < N; ++i)
 159       c[i] = (a[i] * 4) / a[i];
 160
 161 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
 162     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 163 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 164     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 165     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 166     for (int i = 0; i < N; ++i)
 167       d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 168
 169
 170 #pragma acc parallel wait (1) async (1) /* { dg-line l_compute[incr c_compute] } */
 171     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 172 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 173     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 174     /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 175     for (int i = 0; i < N; ++i)
 176       e[i] = a[i] + b[i] + c[i] + d[i];
 177
 178 #pragma acc wait (1)
 179   }
 180
 181   for (int i = 0; i < N; ++i)
 182     {
 183       if (a[i] != 2)
 184         abort ();
 185       if (b[i] != 4)
 186         abort ();
 187       if (c[i] != 4)
 188         abort ();
 189       if (d[i] != 1)
 190         abort ();
 191       if (e[i] != 11)
 192         abort ();
 193     }
 194
 195   for (int i = 0; i < N; ++i)
 196     {
 197       a[i] = 3;
 198       b[i] = 0;
 199     }
 200
 201 #pragma acc data copy (a[0:N]) copy (b[0:N])
 202   {
 203
 204 #pragma acc kernels async
 205 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 206     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 207     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 208     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 209     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 210     for (int i = 0; i < N; ++i)
 211       b[i] = a[i];
 212
 213 #pragma acc wait
 214   }
 215
 216   for (int i = 0; i < N; ++i)
 217     {
 218       if (a[i] != 3)
 219         abort ();
 220       if (b[i] != 3)
 221         abort ();
 222     }
 223
 224   for (int i = 0; i < N; ++i)
 225     {
 226       a[i] = 2;
 227       b[i] = 0;
 228     }
 229
 230 #pragma acc data copy (a[0:N]) copy (b[0:N])
 231   {
 232 #pragma acc kernels async (1)
 233 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 234     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 235     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 236     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 237     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 238     for (int i = 0; i < N; ++i)
 239       b[i] = a[i];
 240
 241 #pragma acc wait (1)
 242   }
 243
 244   for (int i = 0; i < N; ++i)
 245     {
 246       if (a[i] != 2)
 247         abort ();
 248       if (b[i] != 2)
 249         abort ();
 250     }
 251
 252   for (int i = 0; i < N; ++i)
 253     {
 254       a[i] = 3;
 255       b[i] = 0;
 256       c[i] = 0;
 257       d[i] = 0;
 258     }
 259
 260 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
 261   {
 262 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
 263     /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
 264     /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
 265     /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
 266     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
 267        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
 268     /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
 269     for (int i = 0; i < N; ++i)
 270       b[i] = (a[i] * a[i] * a[i]) / a[i];
 271
 272 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
 273     /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
 274     /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
 275     /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
 276     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
 277        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
 278     /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
 279     for (int i = 0; i < N; ++i)
 280       c[i] = (a[i] * 4) / a[i];
 281
 282 #pragma acc kernels async (1)
 283 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 284     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 285     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 286     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 287     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 288     for (int i = 0; i < N; ++i)
 289       d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 290
 291 #pragma acc wait (1)
 292   }
 293
 294   for (int i = 0; i < N; ++i)
 295     {
 296       if (a[i] != 3)
 297         abort ();
 298       if (b[i] != 9)
 299         abort ();
 300       if (c[i] != 4)
 301         abort ();
 302       if (d[i] != 1)
 303         abort ();
 304     }
 305
 306   for (int i = 0; i < N; ++i)
 307     {
 308       a[i] = 2;
 309       b[i] = 0;
 310       c[i] = 0;
 311       d[i] = 0;
 312       e[i] = 0;
 313     }
 314
 315 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
 316   {
 317 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
 318     /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
 319     /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
 320     /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
 321     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
 322        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
 323     /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
 324     for (int i = 0; i < N; ++i)
 325       b[i] = (a[i] * a[i] * a[i]) / a[i];
 326
 327 #pragma acc kernels async (1)
 328 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 329     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 330     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 331     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 332     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 333     for (int i = 0; i < N; ++i)
 334       c[i] = (a[i] * 4) / a[i];
 335
 336 #pragma acc kernels async (1)
 337 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 338     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 339     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 340     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 341     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 342     for (int i = 0; i < N; ++i)
 343       d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 344
 345 #pragma acc kernels wait (1) async (1)
 346 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
 347     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
 348     /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 349     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
 350     /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
 351     for (int i = 0; i < N; ++i)
 352       e[i] = a[i] + b[i] + c[i] + d[i];
 353
 354 #pragma acc wait (1)
 355   }
 356
 357   for (int i = 0; i < N; ++i)
 358     {
 359       if (a[i] != 2)
 360         abort ();
 361       if (b[i] != 4)
 362         abort ();
 363       if (c[i] != 4)
 364         abort ();
 365       if (d[i] != 1)
 366         abort ();
 367       if (e[i] != 11)
 368         abort ();
 369     }
 370
 371   free (a);
 372   free (b);
 373   free (c);
 374   free (d);
 375   free (e);
 376
 377   return 0;
 378 }