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