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

   1 #include <assert.h>
   2 #include <openacc.h>
   3
   4 typedef struct {
   5   int x, y;
   6 } vec2;
   7
   8 typedef struct {
   9   int x, y, z;
  10   int attr[13];
  11 } vec3_attr;
  12
  13
  14 /* Test of gang-private variables declared in local scope with parallel
  15    directive.  */
  16
  17 void local_g_1()
  18 {
  19   int i, arr[32];
  20
  21   for (i = 0; i < 32; i++)
  22     arr[i] = 3;
  23
  24   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
  25   {
  26     int x;
  27
  28     #pragma acc loop gang(static:1)
  29     for (i = 0; i < 32; i++)
  30       x = i * 2;
  31
  32     #pragma acc loop gang(static:1)
  33     for (i = 0; i < 32; i++)
  34       {
  35         if (acc_on_device (acc_device_host))
  36           x = i * 2;
  37         arr[i] += x;
  38       }
  39   }
  40
  41   for (i = 0; i < 32; i++)
  42     assert (arr[i] == 3 + i * 2);
  43 }
  44
  45
  46 /* Test of worker-private variables declared in a local scope, broadcasting
  47    to vector-partitioned mode.  Back-to-back worker loops.  */
  48
  49 void local_w_1()
  50 {
  51   int i, arr[32 * 32 * 32];
  52
  53   for (i = 0; i < 32 * 32 * 32; i++)
  54     arr[i] = i;
  55
  56   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
  57   {
  58     int j;
  59
  60     #pragma acc loop gang
  61     for (i = 0; i < 32; i++)
  62       {
  63         #pragma acc loop worker
  64         for (j = 0; j < 32; j++)
  65           {
  66             int k;
  67             int x = i ^ j * 3;
  68
  69             #pragma acc loop vector
  70             for (k = 0; k < 32; k++)
  71               arr[i * 1024 + j * 32 + k] += x * k;
  72           }
  73
  74         #pragma acc loop worker
  75         for (j = 0; j < 32; j++)
  76           {
  77             int k;
  78             int x = i | j * 5;
  79
  80             #pragma acc loop vector
  81             for (k = 0; k < 32; k++)
  82               arr[i * 1024 + j * 32 + k] += x * k;
  83           }
  84       }
  85   }
  86
  87   for (i = 0; i < 32; i++)
  88     for (int j = 0; j < 32; j++)
  89       for (int k = 0; k < 32; k++)
  90         {
  91           int idx = i * 1024 + j * 32 + k;
  92           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
  93         }
  94 }
  95
  96
  97 /* Test of worker-private variables declared in a local scope, broadcasting
  98    to vector-partitioned mode.  Successive vector loops.  */
  99
 100 void local_w_2()
 101 {
 102   int i, arr[32 * 32 * 32];
 103
 104   for (i = 0; i < 32 * 32 * 32; i++)
 105     arr[i] = i;
 106
 107   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 108   {
 109     int j;
 110
 111     #pragma acc loop gang
 112     for (i = 0; i < 32; i++)
 113       {
 114         #pragma acc loop worker
 115         for (j = 0; j < 32; j++)
 116           {
 117             int k;
 118             int x = i ^ j * 3;
 119
 120             #pragma acc loop vector
 121             for (k = 0; k < 32; k++)
 122               arr[i * 1024 + j * 32 + k] += x * k;
 123
 124             x = i | j * 5;
 125
 126             #pragma acc loop vector
 127             for (k = 0; k < 32; k++)
 128               arr[i * 1024 + j * 32 + k] += x * k;
 129           }
 130       }
 131   }
 132
 133   for (i = 0; i < 32; i++)
 134     for (int j = 0; j < 32; j++)
 135       for (int k = 0; k < 32; k++)
 136         {
 137           int idx = i * 1024 + j * 32 + k;
 138           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 139         }
 140 }
 141
 142
 143 /* Test of worker-private variables declared in a local scope, broadcasting
 144    to vector-partitioned mode.  Aggregate worker variable.  */
 145
 146 void local_w_3()
 147 {
 148   int i, arr[32 * 32 * 32];
 149
 150   for (i = 0; i < 32 * 32 * 32; i++)
 151     arr[i] = i;
 152
 153   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 154   {
 155     int j;
 156
 157     #pragma acc loop gang
 158     for (i = 0; i < 32; i++)
 159       {
 160         #pragma acc loop worker
 161         for (j = 0; j < 32; j++)
 162           {
 163             int k;
 164             vec2 pt;
 165
 166             pt.x = i ^ j * 3;
 167             pt.y = i | j * 5;
 168
 169             #pragma acc loop vector
 170             for (k = 0; k < 32; k++)
 171               arr[i * 1024 + j * 32 + k] += pt.x * k;
 172
 173             #pragma acc loop vector
 174             for (k = 0; k < 32; k++)
 175               arr[i * 1024 + j * 32 + k] += pt.y * k;
 176           }
 177       }
 178   }
 179
 180   for (i = 0; i < 32; i++)
 181     for (int j = 0; j < 32; j++)
 182       for (int k = 0; k < 32; k++)
 183         {
 184           int idx = i * 1024 + j * 32 + k;
 185           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 186         }
 187 }
 188
 189
 190 /* Test of worker-private variables declared in a local scope, broadcasting
 191    to vector-partitioned mode.  Addressable worker variable.  */
 192
 193 void local_w_4()
 194 {
 195   int i, arr[32 * 32 * 32];
 196
 197   for (i = 0; i < 32 * 32 * 32; i++)
 198     arr[i] = i;
 199
 200   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 201   {
 202     int j;
 203
 204     #pragma acc loop gang
 205     for (i = 0; i < 32; i++)
 206       {
 207         #pragma acc loop worker
 208         for (j = 0; j < 32; j++)
 209           {
 210             int k;
 211             vec2 pt, *ptp;
 212
 213             ptp = &pt;
 214
 215             pt.x = i ^ j * 3;
 216
 217             #pragma acc loop vector
 218             for (k = 0; k < 32; k++)
 219               arr[i * 1024 + j * 32 + k] += ptp->x * k;
 220
 221             ptp->y = i | j * 5;
 222
 223             #pragma acc loop vector
 224             for (k = 0; k < 32; k++)
 225               arr[i * 1024 + j * 32 + k] += pt.y * k;
 226           }
 227       }
 228   }
 229
 230   for (i = 0; i < 32; i++)
 231     for (int j = 0; j < 32; j++)
 232       for (int k = 0; k < 32; k++)
 233         {
 234           int idx = i * 1024 + j * 32 + k;
 235           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 236         }
 237 }
 238
 239
 240 /* Test of worker-private variables declared in a local scope, broadcasting
 241    to vector-partitioned mode.  Array worker variable.  */
 242
 243 void local_w_5()
 244 {
 245   int i, arr[32 * 32 * 32];
 246
 247   for (i = 0; i < 32 * 32 * 32; i++)
 248     arr[i] = i;
 249
 250   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 251   {
 252     int j;
 253
 254     #pragma acc loop gang
 255     for (i = 0; i < 32; i++)
 256       {
 257         #pragma acc loop worker
 258         for (j = 0; j < 32; j++)
 259           {
 260             int k;
 261             int pt[2];
 262
 263             pt[0] = i ^ j * 3;
 264
 265             #pragma acc loop vector
 266             for (k = 0; k < 32; k++)
 267               arr[i * 1024 + j * 32 + k] += pt[0] * k;
 268
 269             pt[1] = i | j * 5;
 270
 271             #pragma acc loop vector
 272             for (k = 0; k < 32; k++)
 273               arr[i * 1024 + j * 32 + k] += pt[1] * k;
 274           }
 275       }
 276   }
 277
 278   for (i = 0; i < 32; i++)
 279     for (int j = 0; j < 32; j++)
 280       for (int k = 0; k < 32; k++)
 281         {
 282           int idx = i * 1024 + j * 32 + k;
 283           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 284         }
 285 }
 286
 287
 288 /* Test of gang-private variables declared on loop directive.  */
 289
 290 void loop_g_1()
 291 {
 292   int x = 5, i, arr[32];
 293
 294   for (i = 0; i < 32; i++)
 295     arr[i] = i;
 296
 297   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 298   {
 299     #pragma acc loop gang private(x)
 300     for (i = 0; i < 32; i++)
 301       {
 302         x = i * 2;
 303         arr[i] += x;
 304       }
 305   }
 306
 307   for (i = 0; i < 32; i++)
 308     assert (arr[i] == i * 3);
 309 }
 310
 311
 312 /* Test of gang-private variables declared on loop directive, with broadcasting
 313    to partitioned workers.  */
 314
 315 void loop_g_2()
 316 {
 317   int x = 5, i, arr[32 * 32];
 318
 319   for (i = 0; i < 32 * 32; i++)
 320     arr[i] = i;
 321
 322   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 323   {
 324     #pragma acc loop gang private(x)
 325     for (i = 0; i < 32; i++)
 326       {
 327         x = i * 2;
 328
 329         #pragma acc loop worker
 330         for (int j = 0; j < 32; j++)
 331           arr[i * 32 + j] += x;
 332       }
 333   }
 334
 335   for (i = 0; i < 32 * 32; i++)
 336     assert (arr[i] == i + (i / 32) * 2);
 337 }
 338
 339
 340 /* Test of gang-private variables declared on loop directive, with broadcasting
 341    to partitioned vectors.  */
 342
 343 void loop_g_3()
 344 {
 345   int x = 5, i, arr[32 * 32];
 346
 347   for (i = 0; i < 32 * 32; i++)
 348     arr[i] = i;
 349
 350   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 351   {
 352     #pragma acc loop gang private(x)
 353     for (i = 0; i < 32; i++)
 354       {
 355         x = i * 2;
 356
 357         #pragma acc loop vector
 358         for (int j = 0; j < 32; j++)
 359           arr[i * 32 + j] += x;
 360       }
 361   }
 362
 363   for (i = 0; i < 32 * 32; i++)
 364     assert (arr[i] == i + (i / 32) * 2);
 365 }
 366
 367
 368 /* Test of gang-private addressable variable declared on loop directive, with
 369    broadcasting to partitioned workers.  */
 370
 371 void loop_g_4()
 372 {
 373   int x = 5, i, arr[32 * 32];
 374
 375   for (i = 0; i < 32 * 32; i++)
 376     arr[i] = i;
 377
 378   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 379   {
 380     #pragma acc loop gang private(x)
 381     for (i = 0; i < 32; i++)
 382       {
 383         int *p = &x;
 384
 385         x = i * 2;
 386
 387         #pragma acc loop worker
 388         for (int j = 0; j < 32; j++)
 389           arr[i * 32 + j] += x;
 390
 391         (*p)--;
 392       }
 393   }
 394
 395   for (i = 0; i < 32 * 32; i++)
 396     assert (arr[i] == i + (i / 32) * 2);
 397 }
 398
 399
 400 /* Test of gang-private array variable declared on loop directive, with
 401    broadcasting to partitioned workers.  */
 402
 403 void loop_g_5()
 404 {
 405   int x[8], i, arr[32 * 32];
 406
 407   for (i = 0; i < 32 * 32; i++)
 408     arr[i] = i;
 409
 410   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 411   {
 412     #pragma acc loop gang private(x)
 413     for (i = 0; i < 32; i++)
 414       {
 415         for (int j = 0; j < 8; j++)
 416           x[j] = j * 2;
 417
 418         #pragma acc loop worker
 419         for (int j = 0; j < 32; j++)
 420           arr[i * 32 + j] += x[j % 8];
 421       }
 422   }
 423
 424   for (i = 0; i < 32 * 32; i++)
 425     assert (arr[i] == i + (i % 8) * 2);
 426 }
 427
 428
 429 /* Test of gang-private aggregate variable declared on loop directive, with
 430    broadcasting to partitioned workers.  */
 431
 432 void loop_g_6()
 433 {
 434   int i, arr[32 * 32];
 435   vec3_attr pt;
 436
 437   for (i = 0; i < 32 * 32; i++)
 438     arr[i] = i;
 439
 440   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 441   {
 442     #pragma acc loop gang private(pt)
 443     for (i = 0; i < 32; i++)
 444       {
 445         pt.x = i;
 446         pt.y = i * 2;
 447         pt.z = i * 4;
 448         pt.attr[5] = i * 6;
 449
 450         #pragma acc loop worker
 451         for (int j = 0; j < 32; j++)
 452           arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5];
 453       }
 454   }
 455
 456   for (i = 0; i < 32 * 32; i++)
 457     assert (arr[i] == i + (i / 32) * 13);
 458 }
 459
 460
 461 /* Test of vector-private variables declared on loop directive.  */
 462
 463 void loop_v_1()
 464 {
 465   int x, i, arr[32 * 32 * 32];
 466
 467   for (i = 0; i < 32 * 32 * 32; i++)
 468     arr[i] = i;
 469
 470   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 471   {
 472     int j;
 473
 474     #pragma acc loop gang
 475     for (i = 0; i < 32; i++)
 476       {
 477         #pragma acc loop worker
 478         for (j = 0; j < 32; j++)
 479           {
 480             int k;
 481
 482             #pragma acc loop vector private(x)
 483             for (k = 0; k < 32; k++)
 484               {
 485                 x = i ^ j * 3;
 486                 arr[i * 1024 + j * 32 + k] += x * k;
 487               }
 488
 489             #pragma acc loop vector private(x)
 490             for (k = 0; k < 32; k++)
 491               {
 492                 x = i | j * 5;
 493                 arr[i * 1024 + j * 32 + k] += x * k;
 494               }
 495           }
 496       }
 497   }
 498
 499   for (i = 0; i < 32; i++)
 500     for (int j = 0; j < 32; j++)
 501       for (int k = 0; k < 32; k++)
 502         {
 503           int idx = i * 1024 + j * 32 + k;
 504           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 505         }
 506 }
 507
 508
 509 /* Test of vector-private variables declared on loop directive. Array type.  */
 510
 511 void loop_v_2()
 512 {
 513   int pt[2], i, arr[32 * 32 * 32];
 514
 515   for (i = 0; i < 32 * 32 * 32; i++)
 516     arr[i] = i;
 517
 518   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 519   {
 520     int j;
 521
 522     #pragma acc loop gang
 523     for (i = 0; i < 32; i++)
 524       {
 525         #pragma acc loop worker
 526         for (j = 0; j < 32; j++)
 527           {
 528             int k;
 529
 530             #pragma acc loop vector private(pt)
 531             for (k = 0; k < 32; k++)
 532               {
 533                 pt[0] = i ^ j * 3;
 534                 pt[1] = i | j * 5;
 535                 arr[i * 1024 + j * 32 + k] += pt[0] * k;
 536                 arr[i * 1024 + j * 32 + k] += pt[1] * k;
 537               }
 538           }
 539       }
 540   }
 541
 542   for (i = 0; i < 32; i++)
 543     for (int j = 0; j < 32; j++)
 544       for (int k = 0; k < 32; k++)
 545         {
 546           int idx = i * 1024 + j * 32 + k;
 547           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 548         }
 549 }
 550
 551
 552 /* Test of worker-private variables declared on a loop directive.  */
 553
 554 void loop_w_1()
 555 {
 556   int x = 5, i, arr[32 * 32];
 557
 558   for (i = 0; i < 32 * 32; i++)
 559     arr[i] = i;
 560
 561   #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 562   {
 563     int j;
 564
 565     #pragma acc loop gang
 566     for (i = 0; i < 32; i++)
 567       {
 568         #pragma acc loop worker private(x)
 569         for (j = 0; j < 32; j++)
 570           {
 571             x = i ^ j * 3;
 572             /* Try to ensure 'x' accesses doesn't get optimized into a
 573                temporary.  */
 574             __asm__ __volatile__ ("");
 575             arr[i * 32 + j] += x;
 576           }
 577       }
 578   }
 579
 580   for (i = 0; i < 32 * 32; i++)
 581     assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3));
 582 }
 583
 584
 585 /* Test of worker-private variables declared on a loop directive, broadcasting
 586    to vector-partitioned mode.  */
 587
 588 void loop_w_2()
 589 {
 590   int x = 5, i, arr[32 * 32 * 32];
 591
 592   for (i = 0; i < 32 * 32 * 32; i++)
 593     arr[i] = i;
 594
 595   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 596   {
 597     int j;
 598
 599     #pragma acc loop gang
 600     for (i = 0; i < 32; i++)
 601       {
 602         #pragma acc loop worker private(x)
 603         for (j = 0; j < 32; j++)
 604           {
 605             int k;
 606             x = i ^ j * 3;
 607
 608             #pragma acc loop vector
 609             for (k = 0; k < 32; k++)
 610               arr[i * 1024 + j * 32 + k] += x * k;
 611           }
 612       }
 613   }
 614
 615   for (i = 0; i < 32; i++)
 616     for (int j = 0; j < 32; j++)
 617       for (int k = 0; k < 32; k++)
 618         {
 619           int idx = i * 1024 + j * 32 + k;
 620           assert (arr[idx] == idx + (i ^ j * 3) * k);
 621         }
 622 }
 623
 624
 625 /* Test of worker-private variables declared on a loop directive, broadcasting
 626    to vector-partitioned mode.  Back-to-back worker loops.  */
 627
 628 void loop_w_3()
 629 {
 630   int x = 5, i, arr[32 * 32 * 32];
 631
 632   for (i = 0; i < 32 * 32 * 32; i++)
 633     arr[i] = i;
 634
 635   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 636   {
 637     int j;
 638
 639     #pragma acc loop gang
 640     for (i = 0; i < 32; i++)
 641       {
 642         #pragma acc loop worker private(x)
 643         for (j = 0; j < 32; j++)
 644           {
 645             int k;
 646             x = i ^ j * 3;
 647
 648             #pragma acc loop vector
 649             for (k = 0; k < 32; k++)
 650               arr[i * 1024 + j * 32 + k] += x * k;
 651           }
 652
 653         #pragma acc loop worker private(x)
 654         for (j = 0; j < 32; j++)
 655           {
 656             int k;
 657             x = i | j * 5;
 658
 659             #pragma acc loop vector
 660             for (k = 0; k < 32; k++)
 661               arr[i * 1024 + j * 32 + k] += x * k;
 662           }
 663       }
 664   }
 665
 666   for (i = 0; i < 32; i++)
 667     for (int j = 0; j < 32; j++)
 668       for (int k = 0; k < 32; k++)
 669         {
 670           int idx = i * 1024 + j * 32 + k;
 671           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 672         }
 673 }
 674
 675
 676 /* Test of worker-private variables declared on a loop directive, broadcasting
 677    to vector-partitioned mode.  Successive vector loops.  */
 678
 679 void loop_w_4()
 680 {
 681   int x = 5, i, arr[32 * 32 * 32];
 682
 683   for (i = 0; i < 32 * 32 * 32; i++)
 684     arr[i] = i;
 685
 686   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 687   {
 688     int j;
 689
 690     #pragma acc loop gang
 691     for (i = 0; i < 32; i++)
 692       {
 693         #pragma acc loop worker private(x)
 694         for (j = 0; j < 32; j++)
 695           {
 696             int k;
 697             x = i ^ j * 3;
 698
 699             #pragma acc loop vector
 700             for (k = 0; k < 32; k++)
 701               arr[i * 1024 + j * 32 + k] += x * k;
 702
 703             x = i | j * 5;
 704
 705             #pragma acc loop vector
 706             for (k = 0; k < 32; k++)
 707               arr[i * 1024 + j * 32 + k] += x * k;
 708           }
 709       }
 710   }
 711
 712   for (i = 0; i < 32; i++)
 713     for (int j = 0; j < 32; j++)
 714       for (int k = 0; k < 32; k++)
 715         {
 716           int idx = i * 1024 + j * 32 + k;
 717           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 718         }
 719 }
 720
 721
 722 /* Test of worker-private variables declared on a loop directive, broadcasting
 723    to vector-partitioned mode.  Addressable worker variable.  */
 724
 725 void loop_w_5()
 726 {
 727   int x = 5, i, arr[32 * 32 * 32];
 728
 729   for (i = 0; i < 32 * 32 * 32; i++)
 730     arr[i] = i;
 731
 732   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 733   {
 734     int j;
 735
 736     #pragma acc loop gang
 737     for (i = 0; i < 32; i++)
 738       {
 739         #pragma acc loop worker private(x)
 740         for (j = 0; j < 32; j++)
 741           {
 742             int k;
 743             int *p = &x;
 744
 745             x = i ^ j * 3;
 746
 747             #pragma acc loop vector
 748             for (k = 0; k < 32; k++)
 749               arr[i * 1024 + j * 32 + k] += x * k;
 750
 751             *p = i | j * 5;
 752
 753             #pragma acc loop vector
 754             for (k = 0; k < 32; k++)
 755               arr[i * 1024 + j * 32 + k] += x * k;
 756           }
 757       }
 758   }
 759
 760   for (i = 0; i < 32; i++)
 761     for (int j = 0; j < 32; j++)
 762       for (int k = 0; k < 32; k++)
 763         {
 764           int idx = i * 1024 + j * 32 + k;
 765           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 766         }
 767 }
 768
 769
 770 /* Test of worker-private variables declared on a loop directive, broadcasting
 771    to vector-partitioned mode.  Aggregate worker variable.  */
 772
 773 void loop_w_6()
 774 {
 775   int i, arr[32 * 32 * 32];
 776   vec2 pt;
 777
 778   for (i = 0; i < 32 * 32 * 32; i++)
 779     arr[i] = i;
 780
 781   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 782   {
 783     int j;
 784
 785     #pragma acc loop gang
 786     for (i = 0; i < 32; i++)
 787       {
 788         #pragma acc loop worker private(pt)
 789         for (j = 0; j < 32; j++)
 790           {
 791             int k;
 792
 793             pt.x = i ^ j * 3;
 794             pt.y = i | j * 5;
 795
 796             #pragma acc loop vector
 797             for (k = 0; k < 32; k++)
 798               arr[i * 1024 + j * 32 + k] += pt.x * k;
 799
 800             #pragma acc loop vector
 801             for (k = 0; k < 32; k++)
 802               arr[i * 1024 + j * 32 + k] += pt.y * 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 loop directive, broadcasting
 818    to vector-partitioned mode.  Array worker variable.  */
 819
 820 void loop_w_7()
 821 {
 822   int i, arr[32 * 32 * 32];
 823   int pt[2];
 824
 825   for (i = 0; i < 32 * 32 * 32; i++)
 826     arr[i] = i;
 827
 828   /* "pt" is treated as "present_or_copy" on the parallel directive because it
 829      is an array variable.  */
 830   #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
 831   {
 832     int j;
 833
 834     #pragma acc loop gang
 835     for (i = 0; i < 32; i++)
 836       {
 837         /* But here, it is made private per-worker.  */
 838         #pragma acc loop worker private(pt)
 839         for (j = 0; j < 32; j++)
 840           {
 841             int k;
 842
 843             pt[0] = i ^ j * 3;
 844
 845             #pragma acc loop vector
 846             for (k = 0; k < 32; k++)
 847               arr[i * 1024 + j * 32 + k] += pt[0] * k;
 848
 849             pt[1] = i | j * 5;
 850
 851             #pragma acc loop vector
 852             for (k = 0; k < 32; k++)
 853               arr[i * 1024 + j * 32 + k] += pt[1] * k;
 854           }
 855       }
 856   }
 857
 858   for (i = 0; i < 32; i++)
 859     for (int j = 0; j < 32; j++)
 860       for (int k = 0; k < 32; k++)
 861         {
 862           int idx = i * 1024 + j * 32 + k;
 863           assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
 864         }
 865 }
 866
 867
 868 /* Test of gang-private variables declared on the parallel directive.  */
 869
 870 void parallel_g_1()
 871 {
 872   int x = 5, i, arr[32];
 873
 874   for (i = 0; i < 32; i++)
 875     arr[i] = 3;
 876
 877   #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
 878   {
 879     #pragma acc loop gang(static:1)
 880     for (i = 0; i < 32; i++)
 881       x = i * 2;
 882
 883     #pragma acc loop gang(static:1)
 884     for (i = 0; i < 32; i++)
 885       {
 886         if (acc_on_device (acc_device_host))
 887           x = i * 2;
 888         arr[i] += x;
 889       }
 890   }
 891
 892   for (i = 0; i < 32; i++)
 893     assert (arr[i] == 3 + i * 2);
 894 }
 895
 896
 897 /* Test of gang-private array variable declared on the parallel directive.  */
 898
 899 void parallel_g_2()
 900 {
 901   int x[32], i, arr[32 * 32];
 902
 903   for (i = 0; i < 32 * 32; i++)
 904     arr[i] = i;
 905
 906   #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32)
 907   {
 908     #pragma acc loop gang
 909     for (i = 0; i < 32; i++)
 910       {
 911         int j;
 912         for (j = 0; j < 32; j++)
 913           x[j] = j * 2;
 914
 915         #pragma acc loop worker
 916         for (j = 0; j < 32; j++)
 917           arr[i * 32 + j] += x[31 - j];
 918       }
 919   }
 920
 921   for (i = 0; i < 32 * 32; i++)
 922     assert (arr[i] == i + (31 - (i % 32)) * 2);
 923 }
 924
 925
 926 int main ()
 927 {
 928   local_g_1();
 929   local_w_1();
 930   local_w_2();
 931   local_w_3();
 932   local_w_4();
 933   local_w_5();
 934   loop_g_1();
 935   loop_g_2();
 936   loop_g_3();
 937   loop_g_4();
 938   loop_g_5();
 939   loop_g_6();
 940   loop_v_1();
 941   loop_v_2();
 942   loop_w_1();
 943   loop_w_2();
 944   loop_w_3();
 945   loop_w_4();
 946   loop_w_5();
 947   loop_w_6();
 948   loop_w_7();
 949   parallel_g_1();
 950   parallel_g_2();
 951
 952   return 0;
 953 }