libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c

   1 /* { dg-additional-options "-fopt-info-note-omp" }
   2    { dg-additional-options "--param=openacc-privatization=noisy" }
   3    { dg-additional-options "-foffload=-fopt-info-note-omp" }
   4    { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
   5    for testing/documenting aspects of that functionality.  */
   6
   7 /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
   8    aspects of that functionality.  */
   9
  10 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
  11    passed to 'incr' may be unset, and in that case, it will be set to [...]",
  12    so to maintain compatibility with earlier Tcl releases, we manually
  13    initialize counter variables:
  14    { dg-line l_dummy[variable c_compute 0 c_loop 0] }
  15    { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
  16    "WARNING: dg-line var l_dummy defined, but not used".  */
  17
  18 #include <assert.h>
  19 #include <openacc.h>
  20
  21 typedef struct {
  22   int x, y;
  23 } vec2;
  24
  25 typedef struct {
  26   int x, y, z;
  27   int attr[13];
  28 } vec3_attr;
  29
  30
  31 /* Test of gang-private variables declared in local scope with parallel
  32    directive.  */
  33
  34 void local_g_1()
  35 {
  36   int i, arr[32];
  37
  38   for (i = 0; i < 32; i++)
  39     arr[i] = 3;
  40
  41   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
  42   /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  43   /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
  44   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
  45   {
  46     int x;
  47
  48     #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
  49     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  50     for (i = 0; i < 32; i++)
  51       x = i * 2;
  52
  53     #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
  54     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  55     for (i = 0; i < 32; i++)
  56       {
  57         if (acc_on_device (acc_device_host))
  58           x = i * 2;
  59         arr[i] += x;
  60       }
  61   }
  62
  63   for (i = 0; i < 32; i++)
  64     assert (arr[i] == 3 + i * 2);
  65 }
  66
  67
  68 /* Test of worker-private variables declared in a local scope, broadcasting
  69    to vector-partitioned mode.  Back-to-back worker loops.  */
  70
  71 void local_w_1()
  72 {
  73   int i, arr[32 * 32 * 32];
  74
  75   for (i = 0; i < 32 * 32 * 32; i++)
  76     arr[i] = i;
  77
  78   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
  79   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
  80   {
  81     int j;
  82
  83     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
  84     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  85     for (i = 0; i < 32; i++)
  86       {
  87         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
  88         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  89         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  90         /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  91         for (j = 0; j < 32; j++)
  92           {
  93             int k;
  94             int x = i ^ j * 3;
  95
  96             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
  97             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
  98             for (k = 0; k < 32; k++)
  99               arr[i * 1024 + j * 32 + k] += x * k;
 100           }
 101
 102         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 103         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 104         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 105         /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 106         for (j = 0; j < 32; j++)
 107           {
 108             int k;
 109             int x = i | j * 5;
 110
 111             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 112             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 113             for (k = 0; k < 32; k++)
 114               arr[i * 1024 + j * 32 + k] += x * k;
 115           }
 116       }
 117   }
 118
 119   for (i = 0; i < 32; i++)
 120     for (int j = 0; j < 32; j++)
 121       for (int k = 0; k < 32; k++)
 122         {
 123           int idx = i * 1024 + j * 32 + k;
 124           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 125         }
 126 }
 127
 128
 129 /* Test of worker-private variables declared in a local scope, broadcasting
 130    to vector-partitioned mode.  Successive vector loops.  */
 131
 132 void local_w_2()
 133 {
 134   int i, arr[32 * 32 * 32];
 135
 136   for (i = 0; i < 32 * 32 * 32; i++)
 137     arr[i] = i;
 138
 139   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 140   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 141   {
 142     int j;
 143
 144     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 145     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 146     for (i = 0; i < 32; i++)
 147       {
 148         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 149         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 150         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 151         /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 152         for (j = 0; j < 32; j++)
 153           {
 154             int k;
 155             int x = i ^ j * 3;
 156
 157             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 158             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 159             for (k = 0; k < 32; k++)
 160               arr[i * 1024 + j * 32 + k] += x * k;
 161
 162             x = i | j * 5;
 163
 164             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 165             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 166             for (k = 0; k < 32; k++)
 167               arr[i * 1024 + j * 32 + k] += x * k;
 168           }
 169       }
 170   }
 171
 172   for (i = 0; i < 32; i++)
 173     for (int j = 0; j < 32; j++)
 174       for (int k = 0; k < 32; k++)
 175         {
 176           int idx = i * 1024 + j * 32 + k;
 177           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 178         }
 179 }
 180
 181
 182 /* Test of worker-private variables declared in a local scope, broadcasting
 183    to vector-partitioned mode.  Aggregate worker variable.  */
 184
 185 void local_w_3()
 186 {
 187   int i, arr[32 * 32 * 32];
 188
 189   for (i = 0; i < 32 * 32 * 32; i++)
 190     arr[i] = i;
 191
 192   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 193   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 194   {
 195     int j;
 196
 197     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 198     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 199     for (i = 0; i < 32; i++)
 200       {
 201         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 202         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 203         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 204         /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 205         for (j = 0; j < 32; j++)
 206           {
 207             int k;
 208             vec2 pt;
 209
 210             pt.x = i ^ j * 3;
 211             pt.y = i | j * 5;
 212
 213             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 214             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 215             for (k = 0; k < 32; k++)
 216               arr[i * 1024 + j * 32 + k] += pt.x * k;
 217
 218             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 219             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 220             for (k = 0; k < 32; k++)
 221               arr[i * 1024 + j * 32 + k] += pt.y * k;
 222           }
 223       }
 224   }
 225
 226   for (i = 0; i < 32; i++)
 227     for (int j = 0; j < 32; j++)
 228       for (int k = 0; k < 32; k++)
 229         {
 230           int idx = i * 1024 + j * 32 + k;
 231           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 232         }
 233 }
 234
 235
 236 /* Test of worker-private variables declared in a local scope, broadcasting
 237    to vector-partitioned mode.  Addressable worker variable.  */
 238
 239 void local_w_4()
 240 {
 241   int i, arr[32 * 32 * 32];
 242
 243   for (i = 0; i < 32 * 32 * 32; i++)
 244     arr[i] = i;
 245
 246   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 247   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 248   {
 249     int j;
 250
 251     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 252     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 253     for (i = 0; i < 32; i++)
 254       {
 255         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 256         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 257         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 258         /* { dg-note {variable 'pt' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
 259            { dg-note {variable 'pt' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
 260            { dg-note {variable 'pt' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
 261         /* { dg-note {variable 'ptp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 262         for (j = 0; j < 32; j++)
 263           {
 264             int k;
 265             vec2 pt, *ptp;
 266
 267             ptp = &pt;
 268
 269             pt.x = i ^ j * 3;
 270
 271             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 272             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 273             for (k = 0; k < 32; k++)
 274               arr[i * 1024 + j * 32 + k] += ptp->x * k;
 275
 276             ptp->y = i | j * 5;
 277
 278             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 279             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 280             for (k = 0; k < 32; k++)
 281               arr[i * 1024 + j * 32 + k] += pt.y * k;
 282           }
 283       }
 284   }
 285
 286   for (i = 0; i < 32; i++)
 287     for (int j = 0; j < 32; j++)
 288       for (int k = 0; k < 32; k++)
 289         {
 290           int idx = i * 1024 + j * 32 + k;
 291           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 292         }
 293 }
 294
 295
 296 /* Test of worker-private variables declared in a local scope, broadcasting
 297    to vector-partitioned mode.  Array worker variable.  */
 298
 299 void local_w_5()
 300 {
 301   int i, arr[32 * 32 * 32];
 302
 303   for (i = 0; i < 32 * 32 * 32; i++)
 304     arr[i] = i;
 305
 306   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 307   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 308   {
 309     int j;
 310
 311     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 312     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 313     for (i = 0; i < 32; i++)
 314       {
 315         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 316         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 317         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 318         /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 319         for (j = 0; j < 32; j++)
 320           {
 321             int k;
 322             int pt[2];
 323
 324             pt[0] = i ^ j * 3;
 325
 326             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 327             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 328             for (k = 0; k < 32; k++)
 329               arr[i * 1024 + j * 32 + k] += pt[0] * k;
 330
 331             pt[1] = i | j * 5;
 332
 333             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 334             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 335             for (k = 0; k < 32; k++)
 336               arr[i * 1024 + j * 32 + k] += pt[1] * k;
 337           }
 338       }
 339   }
 340
 341   for (i = 0; i < 32; i++)
 342     for (int j = 0; j < 32; j++)
 343       for (int k = 0; k < 32; k++)
 344         {
 345           int idx = i * 1024 + j * 32 + k;
 346           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 347         }
 348 }
 349
 350
 351 /* Test of gang-private variables declared on loop directive.  */
 352
 353 void loop_g_1()
 354 {
 355   int x = 5, i, arr[32];
 356
 357   for (i = 0; i < 32; i++)
 358     arr[i] = i;
 359
 360   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 361   /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
 362   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 363   {
 364     #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
 365     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 366     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 367     for (i = 0; i < 32; i++)
 368       {
 369         x = i * 2;
 370         arr[i] += x;
 371       }
 372   }
 373
 374   for (i = 0; i < 32; i++)
 375     assert (arr[i] == i * 3);
 376 }
 377
 378
 379 /* Test of gang-private variables declared on loop directive, with broadcasting
 380    to partitioned workers.  */
 381
 382 void loop_g_2()
 383 {
 384   int x = 5, i, arr[32 * 32];
 385
 386   for (i = 0; i < 32 * 32; i++)
 387     arr[i] = i;
 388
 389   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 390   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 391   {
 392     #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
 393     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 394     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 395     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 396     for (i = 0; i < 32; i++)
 397       {
 398         x = i * 2;
 399
 400         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 401         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 402         for (int j = 0; j < 32; j++)
 403           arr[i * 32 + j] += x;
 404       }
 405   }
 406
 407   for (i = 0; i < 32 * 32; i++)
 408     assert (arr[i] == i + (i / 32) * 2);
 409 }
 410
 411
 412 /* Test of gang-private variables declared on loop directive, with broadcasting
 413    to partitioned vectors.  */
 414
 415 void loop_g_3()
 416 {
 417   int x = 5, i, arr[32 * 32];
 418
 419   for (i = 0; i < 32 * 32; i++)
 420     arr[i] = i;
 421
 422   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 423   /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
 424   {
 425     #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
 426     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 427     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 428     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 429     for (i = 0; i < 32; i++)
 430       {
 431         x = i * 2;
 432
 433         #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 434         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 435         for (int j = 0; j < 32; j++)
 436           arr[i * 32 + j] += x;
 437       }
 438   }
 439
 440   for (i = 0; i < 32 * 32; i++)
 441     assert (arr[i] == i + (i / 32) * 2);
 442 }
 443
 444
 445 /* Test of gang-private addressable variable declared on loop directive, with
 446    broadcasting to partitioned workers.  */
 447
 448 void loop_g_4()
 449 {
 450   int x = 5, i, arr[32 * 32];
 451
 452   for (i = 0; i < 32 * 32; i++)
 453     arr[i] = i;
 454
 455   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 456   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 457   {
 458     #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
 459     /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
 460        But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
 461            No longer having address taken: x
 462            Now a gimple register: x
 463        However, 'x' remains in the candidate set:
 464        { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
 465        Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
 466        { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_loop$c_loop }
 467        For nvptx offloading however, we first mark up 'x', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
 468        { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_loop$c_loop }
 469        { dg-bogus {note: variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_loop$c_loop }
 470     */
 471     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 472     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 473     /* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 474     for (i = 0; i < 32; i++)
 475       {
 476         int *p = &x;
 477
 478         x = i * 2;
 479
 480         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 481         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 482         for (int j = 0; j < 32; j++)
 483           arr[i * 32 + j] += x;
 484
 485         (*p)--;
 486       }
 487   }
 488
 489   for (i = 0; i < 32 * 32; i++)
 490     assert (arr[i] == i + (i / 32) * 2);
 491 }
 492
 493
 494 /* Test of gang-private array variable declared on loop directive, with
 495    broadcasting to partitioned workers.  */
 496
 497 void loop_g_5()
 498 {
 499   int x[8], i, arr[32 * 32];
 500
 501   for (i = 0; i < 32 * 32; i++)
 502     arr[i] = i;
 503
 504   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 505   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 506   {
 507     #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
 508     /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
 509        { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
 510        { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } l_loop$c_loop } */
 511     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 512     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 513     for (i = 0; i < 32; i++)
 514       {
 515         for (int j = 0; j < 8; j++)
 516           x[j] = j * 2;
 517
 518         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 519         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 520         for (int j = 0; j < 32; j++)
 521           arr[i * 32 + j] += x[j % 8];
 522       }
 523   }
 524
 525   for (i = 0; i < 32 * 32; i++)
 526     assert (arr[i] == i + (i % 8) * 2);
 527 }
 528
 529
 530 /* Test of gang-private aggregate variable declared on loop directive, with
 531    broadcasting to partitioned workers.  */
 532
 533 void loop_g_6()
 534 {
 535   int i, arr[32 * 32];
 536   vec3_attr pt;
 537
 538   for (i = 0; i < 32 * 32; i++)
 539     arr[i] = i;
 540
 541   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 542   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 543   {
 544     #pragma acc loop gang private(pt) /* { dg-line l_loop[incr c_loop] } */
 545     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 546     /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 547     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 548     for (i = 0; i < 32; i++)
 549       {
 550         pt.x = i;
 551         pt.y = i * 2;
 552         pt.z = i * 4;
 553         pt.attr[5] = i * 6;
 554
 555         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 556         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 557         for (int j = 0; j < 32; j++)
 558           arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5];
 559       }
 560   }
 561
 562   for (i = 0; i < 32 * 32; i++)
 563     assert (arr[i] == i + (i / 32) * 13);
 564 }
 565
 566
 567 /* Test of vector-private variables declared on loop directive.  */
 568
 569 void loop_v_1()
 570 {
 571   int x, i, arr[32 * 32 * 32];
 572
 573   for (i = 0; i < 32 * 32 * 32; i++)
 574     arr[i] = i;
 575
 576   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 577   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 578   {
 579     int j;
 580
 581     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 582     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 583     for (i = 0; i < 32; i++)
 584       {
 585         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 586         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 587         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 588         for (j = 0; j < 32; j++)
 589           {
 590             int k;
 591
 592             #pragma acc loop vector private(x) /* { dg-line l_loop[incr c_loop] } */
 593             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 594             /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 595             for (k = 0; k < 32; k++)
 596               {
 597                 x = i ^ j * 3;
 598                 arr[i * 1024 + j * 32 + k] += x * k;
 599               }
 600
 601             #pragma acc loop vector private(x) /* { dg-line l_loop[incr c_loop] } */
 602             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 603             /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 604             for (k = 0; k < 32; k++)
 605               {
 606                 x = i | j * 5;
 607                 arr[i * 1024 + j * 32 + k] += x * k;
 608               }
 609           }
 610       }
 611   }
 612
 613   for (i = 0; i < 32; i++)
 614     for (int j = 0; j < 32; j++)
 615       for (int k = 0; k < 32; k++)
 616         {
 617           int idx = i * 1024 + j * 32 + k;
 618           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 619         }
 620 }
 621
 622
 623 /* Test of vector-private variables declared on loop directive. Array type.  */
 624
 625 void loop_v_2()
 626 {
 627   int pt[2], i, arr[32 * 32 * 32];
 628
 629   for (i = 0; i < 32 * 32 * 32; i++)
 630     arr[i] = i;
 631
 632   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 633   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 634   {
 635     int j;
 636
 637     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 638     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 639     for (i = 0; i < 32; i++)
 640       {
 641         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
 642         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 643         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 644         for (j = 0; j < 32; j++)
 645           {
 646             int k;
 647
 648             #pragma acc loop vector private(pt) /* { dg-line l_loop[incr c_loop] } */
 649             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 650             /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 651             for (k = 0; k < 32; k++)
 652               {
 653                 pt[0] = i ^ j * 3;
 654                 pt[1] = i | j * 5;
 655                 arr[i * 1024 + j * 32 + k] += pt[0] * k;
 656                 arr[i * 1024 + j * 32 + k] += pt[1] * k;
 657               }
 658           }
 659       }
 660   }
 661
 662   for (i = 0; i < 32; i++)
 663     for (int j = 0; j < 32; j++)
 664       for (int k = 0; k < 32; k++)
 665         {
 666           int idx = i * 1024 + j * 32 + k;
 667           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 668         }
 669 }
 670
 671
 672 /* Test of worker-private variables declared on a loop directive.  */
 673
 674 void loop_w_1()
 675 {
 676   int x = 5, i, arr[32 * 32];
 677
 678   for (i = 0; i < 32 * 32; i++)
 679     arr[i] = i;
 680
 681   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 682   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 683   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
 684   {
 685     int j;
 686
 687     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 688     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 689     for (i = 0; i < 32; i++)
 690       {
 691         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 692         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 693         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 694         for (j = 0; j < 32; j++)
 695           {
 696             x = i ^ j * 3;
 697             /* Try to ensure 'x' accesses doesn't get optimized into a
 698                temporary.  */
 699             __asm__ __volatile__ ("");
 700             arr[i * 32 + j] += x;
 701           }
 702       }
 703   }
 704
 705   for (i = 0; i < 32 * 32; i++)
 706     assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3));
 707 }
 708
 709
 710 /* Test of worker-private variables declared on a loop directive, broadcasting
 711    to vector-partitioned mode.  */
 712
 713 void loop_w_2()
 714 {
 715   int x = 5, i, arr[32 * 32 * 32];
 716
 717   for (i = 0; i < 32 * 32 * 32; i++)
 718     arr[i] = i;
 719
 720   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 721   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 722   {
 723     int j;
 724
 725     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 726     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 727     for (i = 0; i < 32; i++)
 728       {
 729         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 730         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 731         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 732         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 733         for (j = 0; j < 32; j++)
 734           {
 735             int k;
 736             x = i ^ j * 3;
 737
 738             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 739             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 740             for (k = 0; k < 32; k++)
 741               arr[i * 1024 + j * 32 + k] += x * k;
 742           }
 743       }
 744   }
 745
 746   for (i = 0; i < 32; i++)
 747     for (int j = 0; j < 32; j++)
 748       for (int k = 0; k < 32; k++)
 749         {
 750           int idx = i * 1024 + j * 32 + k;
 751           assert (arr[idx] == idx + (i ^ j * 3) * k);
 752         }
 753 }
 754
 755
 756 /* Test of worker-private variables declared on a loop directive, broadcasting
 757    to vector-partitioned mode.  Back-to-back worker loops.  */
 758
 759 void loop_w_3()
 760 {
 761   int x = 5, i, arr[32 * 32 * 32];
 762
 763   for (i = 0; i < 32 * 32 * 32; i++)
 764     arr[i] = i;
 765
 766   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 767   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 768   {
 769     int j;
 770
 771     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 772     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 773     for (i = 0; i < 32; i++)
 774       {
 775         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 776         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 777         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 778         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 779         for (j = 0; j < 32; j++)
 780           {
 781             int k;
 782             x = i ^ j * 3;
 783
 784             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 785             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 786             for (k = 0; k < 32; k++)
 787               arr[i * 1024 + j * 32 + k] += x * k;
 788           }
 789
 790         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 791         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 792         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 793         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 794         for (j = 0; j < 32; j++)
 795           {
 796             int k;
 797             x = i | j * 5;
 798
 799             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 800             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 801             for (k = 0; k < 32; k++)
 802               arr[i * 1024 + j * 32 + k] += x * k;
 803           }
 804       }
 805   }
 806
 807   for (i = 0; i < 32; i++)
 808     for (int j = 0; j < 32; j++)
 809       for (int k = 0; k < 32; k++)
 810         {
 811           int idx = i * 1024 + j * 32 + k;
 812           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 813         }
 814 }
 815
 816
 817 /* Test of worker-private variables declared on a loop directive, broadcasting
 818    to vector-partitioned mode.  Successive vector loops.  */
 819
 820 void loop_w_4()
 821 {
 822   int x = 5, i, arr[32 * 32 * 32];
 823
 824   for (i = 0; i < 32 * 32 * 32; i++)
 825     arr[i] = i;
 826
 827   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 828   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 829   {
 830     int j;
 831
 832     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 833     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 834     for (i = 0; i < 32; i++)
 835       {
 836         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 837         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 838         /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 839         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 840         for (j = 0; j < 32; j++)
 841           {
 842             int k;
 843             x = i ^ j * 3;
 844
 845             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 846             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 847             for (k = 0; k < 32; k++)
 848               arr[i * 1024 + j * 32 + k] += x * k;
 849
 850             x = i | j * 5;
 851
 852             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 853             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 854             for (k = 0; k < 32; k++)
 855               arr[i * 1024 + j * 32 + k] += x * k;
 856           }
 857       }
 858   }
 859
 860   for (i = 0; i < 32; i++)
 861     for (int j = 0; j < 32; j++)
 862       for (int k = 0; k < 32; k++)
 863         {
 864           int idx = i * 1024 + j * 32 + k;
 865           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 866         }
 867 }
 868
 869
 870 /* Test of worker-private variables declared on a loop directive, broadcasting
 871    to vector-partitioned mode.  Addressable worker variable.  */
 872
 873 void loop_w_5()
 874 {
 875   int x = 5, i, arr[32 * 32 * 32];
 876
 877   for (i = 0; i < 32 * 32 * 32; i++)
 878     arr[i] = i;
 879
 880   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 881   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 882   {
 883     int j;
 884
 885     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 886     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 887     for (i = 0; i < 32; i++)
 888       {
 889         #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
 890         /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
 891            { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
 892            { dg-note {variable 'x' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
 893         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 894         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 895         /* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 896         for (j = 0; j < 32; j++)
 897           {
 898             int k;
 899             int *p = &x;
 900
 901             x = i ^ j * 3;
 902
 903             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 904             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 905             for (k = 0; k < 32; k++)
 906               arr[i * 1024 + j * 32 + k] += x * k;
 907
 908             *p = i | j * 5;
 909
 910             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 911             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 912             for (k = 0; k < 32; k++)
 913               arr[i * 1024 + j * 32 + k] += x * k;
 914           }
 915       }
 916   }
 917
 918   for (i = 0; i < 32; i++)
 919     for (int j = 0; j < 32; j++)
 920       for (int k = 0; k < 32; k++)
 921         {
 922           int idx = i * 1024 + j * 32 + k;
 923           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 924         }
 925 }
 926
 927
 928 /* Test of worker-private variables declared on a loop directive, broadcasting
 929    to vector-partitioned mode.  Aggregate worker variable.  */
 930
 931 void loop_w_6()
 932 {
 933   int i, arr[32 * 32 * 32];
 934   vec2 pt;
 935
 936   for (i = 0; i < 32 * 32 * 32; i++)
 937     arr[i] = i;
 938
 939   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 940   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 941   {
 942     int j;
 943
 944     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
 945     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 946     for (i = 0; i < 32; i++)
 947       {
 948         #pragma acc loop worker private(pt) /* { dg-line l_loop[incr c_loop] } */
 949         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 950         /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 951         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 952         for (j = 0; j < 32; j++)
 953           {
 954             int k;
 955
 956             pt.x = i ^ j * 3;
 957             pt.y = i | j * 5;
 958
 959             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 960             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 961             for (k = 0; k < 32; k++)
 962               arr[i * 1024 + j * 32 + k] += pt.x * k;
 963
 964             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
 965             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
 966             for (k = 0; k < 32; k++)
 967               arr[i * 1024 + j * 32 + k] += pt.y * k;
 968           }
 969       }
 970   }
 971
 972   for (i = 0; i < 32; i++)
 973     for (int j = 0; j < 32; j++)
 974       for (int k = 0; k < 32; k++)
 975         {
 976           int idx = i * 1024 + j * 32 + k;
 977           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 978         }
 979 }
 980
 981
 982 /* Test of worker-private variables declared on loop directive, broadcasting
 983    to vector-partitioned mode.  Array worker variable.  */
 984
 985 void loop_w_7()
 986 {
 987   int i, arr[32 * 32 * 32];
 988   int pt[2];
 989
 990   for (i = 0; i < 32 * 32 * 32; i++)
 991     arr[i] = i;
 992
 993   /* "pt" is treated as "present_or_copy" on the parallel directive because it
 994      is an array variable.  */
 995   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
 996   /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
 997   {
 998     int j;
 999
1000     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
1001     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1002     for (i = 0; i < 32; i++)
1003       {
1004         /* But here, it is made private per-worker.  */
1005         #pragma acc loop worker private(pt) /* { dg-line l_loop[incr c_loop] } */
1006         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1007         /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1008         /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1009         for (j = 0; j < 32; j++)
1010           {
1011             int k;
1012
1013             pt[0] = i ^ j * 3;
1014
1015             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
1016             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1017             for (k = 0; k < 32; k++)
1018               arr[i * 1024 + j * 32 + k] += pt[0] * k;
1019
1020             pt[1] = i | j * 5;
1021
1022             #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
1023             /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1024             for (k = 0; k < 32; k++)
1025               arr[i * 1024 + j * 32 + k] += pt[1] * k;
1026           }
1027       }
1028   }
1029
1030   for (i = 0; i < 32; i++)
1031     for (int j = 0; j < 32; j++)
1032       for (int k = 0; k < 32; k++)
1033         {
1034           int idx = i * 1024 + j * 32 + k;
1035           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
1036         }
1037 }
1038
1039
1040 /* Test of gang-private variables declared on the parallel directive.  */
1041
1042 void parallel_g_1()
1043 {
1044   int x = 5, i, arr[32];
1045
1046   for (i = 0; i < 32; i++)
1047     arr[i] = 3;
1048
1049   #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
1050   /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
1051   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
1052   {
1053     #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
1054     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1055     for (i = 0; i < 32; i++)
1056       x = i * 2;
1057
1058     #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
1059     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1060     for (i = 0; i < 32; i++)
1061       {
1062         if (acc_on_device (acc_device_host))
1063           x = i * 2;
1064         arr[i] += x;
1065       }
1066   }
1067
1068   for (i = 0; i < 32; i++)
1069     assert (arr[i] == 3 + i * 2);
1070 }
1071
1072
1073 /* Test of gang-private array variable declared on the parallel directive.  */
1074
1075 void parallel_g_2()
1076 {
1077   int x[32], i, arr[32 * 32];
1078
1079   for (i = 0; i < 32 * 32; i++)
1080     arr[i] = i;
1081
1082   #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
1083   /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
1084   {
1085     #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
1086     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1087     /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1088     for (i = 0; i < 32; i++)
1089       {
1090         int j;
1091         for (j = 0; j < 32; j++)
1092           x[j] = j * 2;
1093
1094         #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
1095         /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1096         for (j = 0; j < 32; j++)
1097           arr[i * 32 + j] += x[31 - j];
1098       }
1099   }
1100
1101   for (i = 0; i < 32 * 32; i++)
1102     assert (arr[i] == i + (31 - (i % 32)) * 2);
1103 }
1104
1105
1106 int main ()
1107 {
1108   local_g_1();
1109   local_w_1();
1110   local_w_2();
1111   local_w_3();
1112   local_w_4();
1113   local_w_5();
1114   loop_g_1();
1115   loop_g_2();
1116   loop_g_3();
1117   loop_g_4();
1118   loop_g_5();
1119   loop_g_6();
1120   loop_v_1();
1121   loop_v_2();
1122   loop_w_1();
1123   loop_w_2();
1124   loop_w_3();
1125   loop_w_4();
1126   loop_w_5();
1127   loop_w_6();
1128   loop_w_7();
1129   parallel_g_1();
1130   parallel_g_2();
1131
1132   return 0;
1133 }