libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c

   1 /* { dg-do run } */
   2 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
   3
   4 #include <openacc.h>
   5 #include <string.h>
   6 #include <stdio.h>
   7 #include <stdlib.h>
   8 #include <stdbool.h>
   9
  10 int
  11 main (int argc, char **argv)
  12 {
  13     int N = 8;
  14     float *a, *b, *c, *d;
  15     int i;
  16
  17     a = (float *) malloc (N * sizeof (float));
  18     b = (float *) malloc (N * sizeof (float));
  19     c = (float *) malloc (N * sizeof (float));
  20
  21     for (i = 0; i < N; i++)
  22     {
  23         a[i] = 3.0;
  24         b[i] = 0.0;
  25     }
  26
  27 #pragma acc data copyin (a[0:N]) copyout (b[0:N])
  28     {
  29 #pragma acc parallel
  30         {
  31             int ii;
  32
  33             for (ii = 0; ii < N; ii++)
  34                 b[ii] = a[ii];
  35         }
  36     }
  37
  38     for (i = 0; i < N; i++)
  39     {
  40         if (b[i] != 3.0)
  41             abort ();
  42     }
  43
  44     if (acc_is_present (&a[0], (N * sizeof (float))))
  45       abort ();
  46
  47     if (acc_is_present (&b[0], (N * sizeof (float))))
  48       abort ();
  49
  50     for (i = 0; i < N; i++)
  51     {
  52         a[i] = 5.0;
  53         b[i] = 1.0;
  54     }
  55
  56 #pragma acc data copyin (a[0:N]) copyout (b[0:N])
  57     {
  58 #pragma acc parallel
  59         {
  60             int ii;
  61
  62             for (ii = 0; ii < N; ii++)
  63                 b[ii] = a[ii];
  64         }
  65     }
  66
  67     for (i = 0; i < N; i++)
  68     {
  69         if (b[i] != 5.0)
  70             abort ();
  71     }
  72
  73     if (acc_is_present (&a[0], (N * sizeof (float))))
  74       abort ();
  75
  76     if (acc_is_present (&b[0], (N * sizeof (float))))
  77       abort ();
  78
  79     for (i = 0; i < N; i++)
  80     {
  81         a[i] = 6.0;
  82         b[i] = 0.0;
  83     }
  84
  85     d = (float *) acc_copyin (&a[0], N * sizeof (float));
  86
  87     for (i = 0; i < N; i++)
  88     {
  89         a[i] = 9.0;
  90     }
  91
  92 #pragma acc data present_or_copyin (a[0:N]) copyout (b[0:N])
  93     {
  94 #pragma acc parallel
  95         {
  96             int ii;
  97
  98             for (ii = 0; ii < N; ii++)
  99                 b[ii] = a[ii];
 100         }
 101     }
 102
 103     for (i = 0; i < N; i++)
 104     {
 105         if (b[i] != 6.0)
 106             abort ();
 107     }
 108
 109     if (!acc_is_present (&a[0], (N * sizeof (float))))
 110       abort ();
 111
 112     if (acc_is_present (&b[0], (N * sizeof (float))))
 113       abort ();
 114
 115     acc_free (d);
 116
 117     for (i = 0; i < N; i++)
 118     {
 119         a[i] = 6.0;
 120         b[i] = 0.0;
 121     }
 122
 123 #pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
 124     {
 125 #pragma acc parallel
 126         {
 127             int ii;
 128
 129             for (ii = 0; ii < N; ii++)
 130                 b[ii] = a[ii];
 131         }
 132     }
 133
 134     for (i = 0; i < N; i++)
 135     {
 136         if (b[i] != 6.0)
 137             abort ();
 138     }
 139
 140     if (acc_is_present (&a[0], (N * sizeof (float))))
 141       abort ();
 142
 143     if (acc_is_present (&b[0], (N * sizeof (float))))
 144       abort ();
 145
 146     for (i = 0; i < N; i++)
 147     {
 148         a[i] = 5.0;
 149         b[i] = 2.0;
 150     }
 151
 152     d = (float *) acc_copyin (&b[0], N * sizeof (float));
 153
 154 #pragma acc data copyin (a[0:N]) present_or_copyout (b[0:N])
 155     {
 156 #pragma acc parallel
 157         {
 158             int ii;
 159
 160             for (ii = 0; ii < N; ii++)
 161                 b[ii] = a[ii];
 162         }
 163     }
 164
 165     for (i = 0; i < N; i++)
 166     {
 167         if (a[i] != 5.0)
 168             abort ();
 169
 170         if (b[i] != 2.0)
 171             abort ();
 172     }
 173
 174     if (acc_is_present (&a[0], (N * sizeof (float))))
 175       abort ();
 176
 177     if (!acc_is_present (&b[0], (N * sizeof (float))))
 178       abort ();
 179
 180     acc_free (d);
 181
 182     if (acc_is_present (&b[0], (N * sizeof (float))))
 183       abort ();
 184
 185     for (i = 0; i < N; i++)
 186     {
 187         a[i] = 3.0;
 188         b[i] = 4.0;
 189     }
 190
 191 #pragma acc data copy (a[0:N]) copyout (b[0:N])
 192     {
 193 #pragma acc parallel
 194         {
 195             int ii;
 196
 197             for (ii = 0; ii < N; ii++)
 198             {
 199                 a[ii] = a[ii] + 1;
 200                 b[ii] = a[ii] + 2;
 201             }
 202         }
 203     }
 204
 205     for (i = 0; i < N; i++)
 206     {
 207         if (a[i] != 4.0)
 208             abort ();
 209
 210         if (b[i] != 6.0)
 211             abort ();
 212     }
 213
 214     if (acc_is_present (&a[0], (N * sizeof (float))))
 215       abort ();
 216
 217     if (acc_is_present (&b[0], (N * sizeof (float))))
 218       abort ();
 219
 220     for (i = 0; i < N; i++)
 221     {
 222         a[i] = 4.0;
 223         b[i] = 7.0;
 224     }
 225
 226 #pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
 227     {
 228 #pragma acc parallel
 229         {
 230             int ii;
 231
 232             for (ii = 0; ii < N; ii++)
 233             {
 234                 a[ii] = a[ii] + 1;
 235                 b[ii] = b[ii] + 2;
 236             }
 237         }
 238     }
 239
 240     for (i = 0; i < N; i++)
 241     {
 242         if (a[i] != 5.0)
 243             abort ();
 244
 245         if (b[i] != 9.0)
 246             abort ();
 247     }
 248
 249     if (acc_is_present (&a[0], (N * sizeof (float))))
 250       abort ();
 251
 252     if (acc_is_present (&b[0], (N * sizeof (float))))
 253       abort ();
 254
 255     for (i = 0; i < N; i++)
 256     {
 257         a[i] = 3.0;
 258         b[i] = 7.0;
 259     }
 260
 261     d = (float *) acc_copyin (&a[0], N * sizeof (float));
 262     d = (float *) acc_copyin (&b[0], N * sizeof (float));
 263
 264 #pragma acc data present_or_copy (a[0:N]) present_or_copy (b[0:N])
 265     {
 266 #pragma acc parallel
 267         {
 268             int ii;
 269
 270             for (ii = 0; ii < N; ii++)
 271             {
 272                 a[ii] = a[ii] + 1;
 273                 b[ii] = b[ii] + 2;
 274             }
 275         }
 276     }
 277
 278     for (i = 0; i < N; i++)
 279     {
 280         if (a[i] != 3.0)
 281             abort ();
 282
 283         if (b[i] != 7.0)
 284             abort ();
 285     }
 286
 287     if (!acc_is_present (&a[0], (N * sizeof (float))))
 288       abort ();
 289
 290     if (!acc_is_present (&b[0], (N * sizeof (float))))
 291       abort ();
 292
 293     d = (float *) acc_deviceptr (&a[0]);
 294     acc_unmap_data (&a[0]);
 295     acc_free (d);
 296
 297     d = (float *) acc_deviceptr (&b[0]);
 298     acc_unmap_data (&b[0]);
 299     acc_free (d);
 300
 301
 302     for (i = 0; i < N; i++)
 303     {
 304         a[i] = 3.0;
 305         b[i] = 7.0;
 306     }
 307
 308 #pragma acc data copyin (a[0:N]) create (c[0:N]) copyout (b[0:N])
 309     {
 310 #pragma acc parallel
 311         {
 312             int ii;
 313
 314             for (ii = 0; ii < N; ii++)
 315             {
 316                 c[ii] = a[ii];
 317                 b[ii] = c[ii];
 318             }
 319         }
 320     }
 321
 322     for (i = 0; i < N; i++)
 323     {
 324         if (a[i] != 3.0)
 325             abort ();
 326
 327         if (b[i] != 3.0)
 328             abort ();
 329     }
 330
 331     if (acc_is_present (&a[0], (N * sizeof (float))))
 332       abort ();
 333
 334     if (acc_is_present (&b[0], (N * sizeof (float))))
 335       abort ();
 336
 337     if (acc_is_present (&c[0], (N * sizeof (float))))
 338       abort ();
 339
 340     for (i = 0; i < N; i++)
 341     {
 342         a[i] = 4.0;
 343         b[i] = 8.0;
 344     }
 345
 346 #pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
 347     {
 348 #pragma acc parallel
 349         {
 350             int ii;
 351
 352             for (ii = 0; ii < N; ii++)
 353             {
 354                 c[ii] = a[ii];
 355                 b[ii] = c[ii];
 356             }
 357         }
 358     }
 359
 360     for (i = 0; i < N; i++)
 361     {
 362         if (a[i] != 4.0)
 363             abort ();
 364
 365         if (b[i] != 4.0)
 366             abort ();
 367     }
 368
 369     if (acc_is_present (&a[0], (N * sizeof (float))))
 370       abort ();
 371
 372     if (acc_is_present (&b[0], (N * sizeof (float))))
 373       abort ();
 374
 375     if (acc_is_present (&c[0], (N * sizeof (float))))
 376       abort ();
 377
 378     for (i = 0; i < N; i++)
 379     {
 380         a[i] = 2.0;
 381         b[i] = 5.0;
 382     }
 383
 384     d = (float *) acc_malloc (N * sizeof (float));
 385     acc_map_data (c, d, N * sizeof (float));
 386
 387 #pragma acc data copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
 388     {
 389 #pragma acc parallel
 390         {
 391             int ii;
 392
 393             for (ii = 0; ii < N; ii++)
 394             {
 395                 c[ii] = a[ii];
 396                 b[ii] = c[ii];
 397             }
 398         }
 399     }
 400
 401     for (i = 0; i < N; i++)
 402     {
 403         if (a[i] != 2.0)
 404             abort ();
 405
 406         if (b[i] != 2.0)
 407             abort ();
 408     }
 409
 410     if (acc_is_present (a, (N * sizeof (float))))
 411       abort ();
 412
 413     if (acc_is_present (b, (N * sizeof (float))))
 414       abort ();
 415
 416     if (!acc_is_present (c, (N * sizeof (float))))
 417       abort ();
 418
 419     d = (float *) acc_deviceptr (c);
 420
 421     acc_unmap_data (c);
 422
 423     acc_free (d);
 424
 425     for (i = 0; i < N; i++)
 426     {
 427         a[i] = 4.0;
 428         b[i] = 8.0;
 429     }
 430
 431     d = (float *) acc_malloc (N * sizeof (float));
 432     acc_map_data (c, d, N * sizeof (float));
 433
 434 #pragma acc data copyin (a[0:N]) present (c[0:N]) copyout (b[0:N])
 435     {
 436 #pragma acc parallel
 437         {
 438             int ii;
 439
 440             for (ii = 0; ii < N; ii++)
 441             {
 442                 c[ii] = a[ii];
 443                 b[ii] = c[ii];
 444             }
 445         }
 446     }
 447
 448     for (i = 0; i < N; i++)
 449     {
 450         if (a[i] != 4.0)
 451             abort ();
 452
 453         if (b[i] != 4.0)
 454             abort ();
 455     }
 456
 457     if (acc_is_present (a, (N * sizeof (float))))
 458       abort ();
 459
 460     if (acc_is_present (b, (N * sizeof (float))))
 461       abort ();
 462
 463     if (!acc_is_present (c, (N * sizeof (float))))
 464       abort ();
 465
 466     acc_unmap_data (c);
 467
 468     if (acc_is_present (c, (N * sizeof (float))))
 469       abort ();
 470
 471     acc_free (d);
 472
 473     d = (float *) acc_malloc (N * sizeof (float));
 474     acc_map_data (c, d, N * sizeof (float));
 475
 476     if (!acc_is_present (c, (N * sizeof (float))))
 477       abort ();
 478
 479     d = (float *) acc_malloc (N * sizeof (float));
 480     acc_map_data (b, d, N * sizeof (float));
 481
 482     if (!acc_is_present (b, (N * sizeof (float))))
 483       abort ();
 484
 485     d = (float *) acc_malloc (N * sizeof (float));
 486     acc_map_data (a, d, N * sizeof (float));
 487
 488     if (!acc_is_present (a, (N * sizeof (float))))
 489       abort ();
 490
 491 #pragma acc data present (a[0:N]) present (c[0:N]) present (b[0:N])
 492     {
 493 #pragma acc parallel
 494         {
 495             int ii;
 496
 497             for (ii = 0; ii < N; ii++)
 498             {
 499                 a[ii] = 1.0;
 500                 c[ii] = 2.0;
 501                 b[ii] = 4.0;
 502             }
 503         }
 504     }
 505
 506     if (!acc_is_present (a, (N * sizeof (float))))
 507       abort ();
 508
 509     if (!acc_is_present (b, (N * sizeof (float))))
 510       abort ();
 511
 512     if (!acc_is_present (c, (N * sizeof (float))))
 513       abort ();
 514
 515     acc_copyout (b, N * sizeof (float));
 516
 517     for (i = 0; i < N; i++)
 518     {
 519         if (a[i] != 4.0)
 520             abort ();
 521
 522         if (b[i] != 4.0)
 523             abort ();
 524     }
 525
 526     d = (float *) acc_deviceptr (a);
 527
 528     acc_unmap_data (a);
 529
 530     acc_free (d);
 531
 532     d = (float *) acc_deviceptr (c);
 533
 534     acc_unmap_data (c);
 535
 536     acc_free (d);
 537
 538     for (i = 0; i < N; i++)
 539     {
 540         a[i] = 3.0;
 541         b[i] = 6.0;
 542     }
 543
 544     d = (float *) acc_malloc (N * sizeof (float));
 545
 546 #pragma acc parallel copyin (a[0:N]) deviceptr (d) copyout (b[0:N])
 547     {
 548         int ii;
 549
 550         for (ii = 0; ii < N; ii++)
 551         {
 552             d[ii] = a[ii];
 553             b[ii] = d[ii];
 554         }
 555     }
 556
 557     for (i = 0; i < N; i++)
 558     {
 559         if (a[i] != 3.0)
 560             abort ();
 561
 562         if (b[i] != 3.0)
 563             abort ();
 564     }
 565
 566     if (acc_is_present (a, (N * sizeof (float))))
 567       abort ();
 568
 569     if (acc_is_present (b, (N * sizeof (float))))
 570       abort ();
 571
 572     acc_free (d);
 573
 574     for (i = 0; i < N; i++)
 575     {
 576         a[i] = 6.0;
 577         b[i] = 0.0;
 578     }
 579
 580     d = (float *) acc_copyin (&a[0], N * sizeof (float));
 581
 582     for (i = 0; i < N; i++)
 583     {
 584         a[i] = 9.0;
 585     }
 586
 587 #pragma acc data pcopyin (a[0:N]) copyout (b[0:N])
 588     {
 589 #pragma acc parallel
 590         {
 591             int ii;
 592
 593             for (ii = 0; ii < N; ii++)
 594                 b[ii] = a[ii];
 595         }
 596     }
 597
 598     for (i = 0; i < N; i++)
 599     {
 600         if (b[i] != 6.0)
 601             abort ();
 602     }
 603
 604     if (!acc_is_present (&a[0], (N * sizeof (float))))
 605       abort ();
 606
 607     if (acc_is_present (&b[0], (N * sizeof (float))))
 608       abort ();
 609
 610     acc_free (d);
 611
 612     for (i = 0; i < N; i++)
 613     {
 614         a[i] = 6.0;
 615         b[i] = 0.0;
 616     }
 617
 618 #pragma acc data copyin (a[0:N]) pcopyout (b[0:N])
 619     {
 620 #pragma acc parallel
 621         {
 622             int ii;
 623
 624             for (ii = 0; ii < N; ii++)
 625                 b[ii] = a[ii];
 626         }
 627     }
 628
 629     for (i = 0; i < N; i++)
 630     {
 631         if (b[i] != 6.0)
 632             abort ();
 633     }
 634
 635     if (acc_is_present (&a[0], (N * sizeof (float))))
 636       abort ();
 637
 638     if (acc_is_present (&b[0], (N * sizeof (float))))
 639       abort ();
 640
 641     for (i = 0; i < N; i++)
 642     {
 643         a[i] = 5.0;
 644         b[i] = 7.0;
 645     }
 646
 647 #pragma acc data copyin (a[0:N]) pcreate (c[0:N]) copyout (b[0:N])
 648     {
 649 #pragma acc parallel
 650         {
 651             int ii;
 652
 653             for (ii = 0; ii < N; ii++)
 654             {
 655                 c[ii] = a[ii];
 656                 b[ii] = c[ii];
 657             }
 658         }
 659     }
 660
 661     for (i = 0; i < N; i++)
 662     {
 663         if (a[i] != 5.0)
 664             abort ();
 665
 666         if (b[i] != 5.0)
 667             abort ();
 668     }
 669
 670     if (acc_is_present (&a[0], (N * sizeof (float))))
 671       abort ();
 672
 673     if (acc_is_present (&b[0], (N * sizeof (float))))
 674       abort ();
 675
 676     if (acc_is_present (&c[0], (N * sizeof (float))))
 677       abort ();
 678
 679     return 0;
 680 }