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

   1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
   2 /* { dg-additional-options "-lcuda" } */
   3
   4 #include <openacc.h>
   5 #include <stdlib.h>
   6 #include "cuda.h"
   7
   8 #include <stdio.h>
   9 #include <sys/time.h>
  10
  11 int
  12 main (int argc, char **argv)
  13 {
  14     CUresult r;
  15     CUstream stream1;
  16     int N = 128; //1024 * 1024;
  17     float *a, *b, *c, *d, *e;
  18     int i;
  19     int nbytes;
  20
  21     acc_init (acc_device_nvidia);
  22
  23     nbytes = N * sizeof (float);
  24
  25     a = (float *) malloc (nbytes);
  26     b = (float *) malloc (nbytes);
  27     c = (float *) malloc (nbytes);
  28     d = (float *) malloc (nbytes);
  29     e = (float *) malloc (nbytes);
  30
  31     for (i = 0; i < N; i++)
  32     {
  33         a[i] = 3.0;
  34         b[i] = 0.0;
  35     }
  36
  37 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
  38     {
  39
  40 #pragma acc parallel async
  41     {
  42         int ii;
  43
  44         for (ii = 0; ii < N; ii++)
  45             b[ii] = a[ii];
  46     }
  47
  48 #pragma acc wait
  49
  50     }
  51
  52     for (i = 0; i < N; i++)
  53     {
  54         if (a[i] != 3.0)
  55             abort ();
  56
  57         if (b[i] != 3.0)
  58             abort ();
  59     }
  60
  61     for (i = 0; i < N; i++)
  62     {
  63         a[i] = 2.0;
  64         b[i] = 0.0;
  65     }
  66
  67 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
  68     {
  69
  70 #pragma acc parallel async (1)
  71     {
  72         int ii;
  73
  74         for (ii = 0; ii < N; ii++)
  75             b[ii] = a[ii];
  76     }
  77
  78 #pragma acc wait (1)
  79
  80     }
  81
  82     for (i = 0; i < N; i++)
  83     {
  84         if (a[i] != 2.0)
  85             abort ();
  86
  87         if (b[i] != 2.0)
  88             abort ();
  89     }
  90
  91     for (i = 0; i < N; i++)
  92     {
  93         a[i] = 3.0;
  94         b[i] = 0.0;
  95         c[i] = 0.0;
  96         d[i] = 0.0;
  97     }
  98
  99 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
 100     {
 101
 102 #pragma acc parallel async (1)
 103     {
 104         int ii;
 105
 106         for (ii = 0; ii < N; ii++)
 107             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 108     }
 109
 110 #pragma acc parallel async (1)
 111     {
 112         int ii;
 113
 114         for (ii = 0; ii < N; ii++)
 115             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 116     }
 117
 118
 119 #pragma acc parallel async (1)
 120     {
 121         int ii;
 122
 123         for (ii = 0; ii < N; ii++)
 124             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 125     }
 126
 127 #pragma acc wait (1)
 128
 129     }
 130
 131     for (i = 0; i < N; i++)
 132     {
 133         if (a[i] != 3.0)
 134             abort ();
 135
 136         if (b[i] != 9.0)
 137             abort ();
 138
 139         if (c[i] != 4.0)
 140             abort ();
 141
 142         if (d[i] != 1.0)
 143             abort ();
 144     }
 145
 146     for (i = 0; i < N; i++)
 147     {
 148         a[i] = 2.0;
 149         b[i] = 0.0;
 150         c[i] = 0.0;
 151         d[i] = 0.0;
 152         e[i] = 0.0;
 153     }
 154
 155 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
 156     {
 157
 158 #pragma acc parallel async (1)
 159     {
 160         int ii;
 161
 162         for (ii = 0; ii < N; ii++)
 163             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 164     }
 165
 166 #pragma acc parallel async (1)
 167     {
 168         int ii;
 169
 170         for (ii = 0; ii < N; ii++)
 171             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 172     }
 173
 174 #pragma acc parallel async (1)
 175     {
 176         int ii;
 177
 178         for (ii = 0; ii < N; ii++)
 179             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 180     }
 181
 182 #pragma acc parallel wait (1) async (1)
 183     {
 184         int ii;
 185
 186         for (ii = 0; ii < N; ii++)
 187             e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 188     }
 189
 190 #pragma acc wait (1)
 191
 192     }
 193
 194     for (i = 0; i < N; i++)
 195     {
 196         if (a[i] != 2.0)
 197             abort ();
 198
 199         if (b[i] != 4.0)
 200             abort ();
 201
 202         if (c[i] != 4.0)
 203             abort ();
 204
 205         if (d[i] != 1.0)
 206             abort ();
 207
 208         if (e[i] != 11.0)
 209             abort ();
 210     }
 211
 212
 213     r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
 214     if (r != CUDA_SUCCESS)
 215     {
 216         fprintf (stderr, "cuStreamCreate failed: %d\n", r);
 217         abort ();
 218     }
 219
 220     acc_set_cuda_stream (1, stream1);
 221
 222     for (i = 0; i < N; i++)
 223     {
 224         a[i] = 5.0;
 225         b[i] = 0.0;
 226     }
 227
 228 #pragma acc data copy (a[0:N], b[0:N]) copyin (N)
 229     {
 230
 231 #pragma acc parallel async (1)
 232     {
 233         int ii;
 234
 235         for (ii = 0; ii < N; ii++)
 236             b[ii] = a[ii];
 237     }
 238
 239 #pragma acc wait (1)
 240
 241     }
 242
 243     for (i = 0; i < N; i++)
 244     {
 245         if (a[i] != 5.0)
 246             abort ();
 247
 248         if (b[i] != 5.0)
 249             abort ();
 250     }
 251
 252     for (i = 0; i < N; i++)
 253     {
 254         a[i] = 7.0;
 255         b[i] = 0.0;
 256         c[i] = 0.0;
 257         d[i] = 0.0;
 258     }
 259
 260 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
 261     {
 262
 263 #pragma acc parallel async (1)
 264     {
 265         int ii;
 266
 267         for (ii = 0; ii < N; ii++)
 268             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 269     }
 270
 271 #pragma acc parallel async (1)
 272     {
 273         int ii;
 274
 275         for (ii = 0; ii < N; ii++)
 276             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 277     }
 278
 279 #pragma acc parallel async (1)
 280     {
 281         int ii;
 282
 283         for (ii = 0; ii < N; ii++)
 284             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 285     }
 286
 287 #pragma acc wait (1)
 288
 289     }
 290
 291     for (i = 0; i < N; i++)
 292     {
 293         if (a[i] != 7.0)
 294             abort ();
 295
 296         if (b[i] != 49.0)
 297             abort ();
 298
 299         if (c[i] != 4.0)
 300             abort ();
 301
 302         if (d[i] != 1.0)
 303             abort ();
 304     }
 305
 306     for (i = 0; i < N; i++)
 307     {
 308         a[i] = 3.0;
 309         b[i] = 0.0;
 310         c[i] = 0.0;
 311         d[i] = 0.0;
 312         e[i] = 0.0;
 313     }
 314
 315 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
 316     {
 317
 318 #pragma acc parallel async (1)
 319     {
 320         int ii;
 321
 322         for (ii = 0; ii < N; ii++)
 323             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 324     }
 325
 326 #pragma acc parallel async (1)
 327     {
 328         int ii;
 329
 330         for (ii = 0; ii < N; ii++)
 331             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 332     }
 333
 334 #pragma acc parallel async (1)
 335     {
 336         int ii;
 337
 338         for (ii = 0; ii < N; ii++)
 339             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 340     }
 341
 342 #pragma acc parallel wait (1) async (1)
 343     {
 344         int ii;
 345
 346         for (ii = 0; ii < N; ii++)
 347             e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 348     }
 349
 350 #pragma acc wait (1)
 351
 352     }
 353
 354     for (i = 0; i < N; i++)
 355     {
 356         if (a[i] != 3.0)
 357             abort ();
 358
 359         if (b[i] != 9.0)
 360             abort ();
 361
 362         if (c[i] != 4.0)
 363             abort ();
 364
 365         if (d[i] != 1.0)
 366             abort ();
 367
 368         if (e[i] != 17.0)
 369             abort ();
 370     }
 371
 372     for (i = 0; i < N; i++)
 373     {
 374         a[i] = 4.0;
 375         b[i] = 0.0;
 376         c[i] = 0.0;
 377         d[i] = 0.0;
 378         e[i] = 0.0;
 379     }
 380
 381 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
 382     {
 383
 384 #pragma acc parallel async (1)
 385     {
 386         int ii;
 387
 388         for (ii = 0; ii < N; ii++)
 389             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 390     }
 391
 392 #pragma acc parallel async (1)
 393     {
 394         int ii;
 395
 396         for (ii = 0; ii < N; ii++)
 397             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 398     }
 399
 400 #pragma acc update host (a[0:N], b[0:N], c[0:N]) wait (1)
 401
 402     }
 403
 404     for (i = 0; i < N; i++)
 405     {
 406         if (a[i] != 4.0)
 407             abort ();
 408
 409         if (b[i] != 16.0)
 410             abort ();
 411
 412         if (c[i] != 4.0)
 413             abort ();
 414     }
 415
 416
 417     for (i = 0; i < N; i++)
 418     {
 419         a[i] = 5.0;
 420         b[i] = 0.0;
 421         c[i] = 0.0;
 422         d[i] = 0.0;
 423         e[i] = 0.0;
 424     }
 425
 426 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
 427     {
 428
 429 #pragma acc parallel async (1)
 430     {
 431         int ii;
 432
 433         for (ii = 0; ii < N; ii++)
 434             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 435     }
 436
 437 #pragma acc parallel async (1)
 438     {
 439         int ii;
 440
 441         for (ii = 0; ii < N; ii++)
 442             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 443     }
 444
 445 #pragma acc update host (a[0:N], b[0:N], c[0:N]) async (1)
 446
 447 #pragma acc wait (1)
 448
 449     }
 450
 451     for (i = 0; i < N; i++)
 452     {
 453         if (a[i] != 5.0)
 454             abort ();
 455
 456         if (b[i] != 25.0)
 457             abort ();
 458
 459         if (c[i] != 4.0)
 460             abort ();
 461     }
 462
 463     for (i = 0; i < N; i++)
 464     {
 465         a[i] = 3.0;
 466         b[i] = 0.0;
 467     }
 468
 469 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
 470     {
 471
 472 #pragma acc kernels async
 473     {
 474         int ii;
 475
 476         for (ii = 0; ii < N; ii++)
 477             b[ii] = a[ii];
 478     }
 479
 480 #pragma acc wait
 481
 482     }
 483
 484     for (i = 0; i < N; i++)
 485     {
 486         if (a[i] != 3.0)
 487             abort ();
 488
 489         if (b[i] != 3.0)
 490             abort ();
 491     }
 492
 493     for (i = 0; i < N; i++)
 494     {
 495         a[i] = 2.0;
 496         b[i] = 0.0;
 497     }
 498
 499 #pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
 500     {
 501
 502 #pragma acc kernels async (1)
 503     {
 504         int ii;
 505
 506         for (ii = 0; ii < N; ii++)
 507             b[ii] = a[ii];
 508     }
 509
 510 #pragma acc wait (1)
 511
 512     }
 513
 514     for (i = 0; i < N; i++)
 515     {
 516         if (a[i] != 2.0)
 517             abort ();
 518
 519         if (b[i] != 2.0)
 520             abort ();
 521     }
 522
 523     for (i = 0; i < N; i++)
 524     {
 525         a[i] = 3.0;
 526         b[i] = 0.0;
 527         c[i] = 0.0;
 528         d[i] = 0.0;
 529     }
 530
 531 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
 532     {
 533
 534 #pragma acc kernels async (1)
 535     {
 536         int ii;
 537
 538         for (ii = 0; ii < N; ii++)
 539             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 540     }
 541
 542 #pragma acc kernels async (1)
 543     {
 544         int ii;
 545
 546         for (ii = 0; ii < N; ii++)
 547             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 548     }
 549
 550
 551 #pragma acc kernels async (1)
 552     {
 553         int ii;
 554
 555         for (ii = 0; ii < N; ii++)
 556             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 557     }
 558
 559 #pragma acc wait (1)
 560
 561     }
 562
 563     for (i = 0; i < N; i++)
 564     {
 565         if (a[i] != 3.0)
 566             abort ();
 567
 568         if (b[i] != 9.0)
 569             abort ();
 570
 571         if (c[i] != 4.0)
 572             abort ();
 573
 574         if (d[i] != 1.0)
 575             abort ();
 576     }
 577
 578     for (i = 0; i < N; i++)
 579     {
 580         a[i] = 2.0;
 581         b[i] = 0.0;
 582         c[i] = 0.0;
 583         d[i] = 0.0;
 584         e[i] = 0.0;
 585     }
 586
 587 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
 588     {
 589
 590 #pragma acc kernels async (1)
 591     {
 592         int ii;
 593
 594         for (ii = 0; ii < N; ii++)
 595             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 596     }
 597
 598 #pragma acc kernels async (1)
 599     {
 600         int ii;
 601
 602         for (ii = 0; ii < N; ii++)
 603             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 604     }
 605
 606 #pragma acc kernels async (1)
 607     {
 608         int ii;
 609
 610         for (ii = 0; ii < N; ii++)
 611             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 612     }
 613
 614 #pragma acc kernels wait (1) async (1)
 615     {
 616         int ii;
 617
 618         for (ii = 0; ii < N; ii++)
 619             e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 620     }
 621
 622 #pragma acc wait (1)
 623
 624     }
 625
 626     for (i = 0; i < N; i++)
 627     {
 628         if (a[i] != 2.0)
 629             abort ();
 630
 631         if (b[i] != 4.0)
 632             abort ();
 633
 634         if (c[i] != 4.0)
 635             abort ();
 636
 637         if (d[i] != 1.0)
 638             abort ();
 639
 640         if (e[i] != 11.0)
 641             abort ();
 642     }
 643
 644
 645     r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
 646     if (r != CUDA_SUCCESS)
 647     {
 648         fprintf (stderr, "cuStreamCreate failed: %d\n", r);
 649         abort ();
 650     }
 651
 652     acc_set_cuda_stream (1, stream1);
 653
 654     for (i = 0; i < N; i++)
 655     {
 656         a[i] = 5.0;
 657         b[i] = 0.0;
 658     }
 659
 660 #pragma acc data copy (a[0:N], b[0:N]) copyin (N)
 661     {
 662
 663 #pragma acc kernels async (1)
 664     {
 665         int ii;
 666
 667         for (ii = 0; ii < N; ii++)
 668             b[ii] = a[ii];
 669     }
 670
 671 #pragma acc wait (1)
 672
 673     }
 674
 675     for (i = 0; i < N; i++)
 676     {
 677         if (a[i] != 5.0)
 678             abort ();
 679
 680         if (b[i] != 5.0)
 681             abort ();
 682     }
 683
 684     for (i = 0; i < N; i++)
 685     {
 686         a[i] = 7.0;
 687         b[i] = 0.0;
 688         c[i] = 0.0;
 689         d[i] = 0.0;
 690     }
 691
 692 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
 693     {
 694
 695 #pragma acc kernels async (1)
 696     {
 697         int ii;
 698
 699         for (ii = 0; ii < N; ii++)
 700             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 701     }
 702
 703 #pragma acc kernels async (1)
 704     {
 705         int ii;
 706
 707         for (ii = 0; ii < N; ii++)
 708             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 709     }
 710
 711 #pragma acc kernels async (1)
 712     {
 713         int ii;
 714
 715         for (ii = 0; ii < N; ii++)
 716             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 717     }
 718
 719 #pragma acc wait (1)
 720
 721     }
 722
 723     for (i = 0; i < N; i++)
 724     {
 725         if (a[i] != 7.0)
 726             abort ();
 727
 728         if (b[i] != 49.0)
 729             abort ();
 730
 731         if (c[i] != 4.0)
 732             abort ();
 733
 734         if (d[i] != 1.0)
 735             abort ();
 736     }
 737
 738     for (i = 0; i < N; i++)
 739     {
 740         a[i] = 3.0;
 741         b[i] = 0.0;
 742         c[i] = 0.0;
 743         d[i] = 0.0;
 744         e[i] = 0.0;
 745     }
 746
 747 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
 748     {
 749
 750 #pragma acc kernels async (1)
 751     {
 752         int ii;
 753
 754         for (ii = 0; ii < N; ii++)
 755             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 756     }
 757
 758 #pragma acc kernels async (1)
 759     {
 760         int ii;
 761
 762         for (ii = 0; ii < N; ii++)
 763             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 764     }
 765
 766 #pragma acc kernels async (1)
 767     {
 768         int ii;
 769
 770         for (ii = 0; ii < N; ii++)
 771             d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
 772     }
 773
 774 #pragma acc kernels wait (1) async (1)
 775     {
 776         int ii;
 777
 778         for (ii = 0; ii < N; ii++)
 779             e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
 780     }
 781
 782 #pragma acc wait (1)
 783
 784     }
 785
 786     for (i = 0; i < N; i++)
 787     {
 788         if (a[i] != 3.0)
 789             abort ();
 790
 791         if (b[i] != 9.0)
 792             abort ();
 793
 794         if (c[i] != 4.0)
 795             abort ();
 796
 797         if (d[i] != 1.0)
 798             abort ();
 799
 800         if (e[i] != 17.0)
 801             abort ();
 802     }
 803
 804     for (i = 0; i < N; i++)
 805     {
 806         a[i] = 4.0;
 807         b[i] = 0.0;
 808         c[i] = 0.0;
 809         d[i] = 0.0;
 810         e[i] = 0.0;
 811     }
 812
 813 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
 814     {
 815
 816 #pragma acc kernels async (1)
 817     {
 818         int ii;
 819
 820         for (ii = 0; ii < N; ii++)
 821             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 822     }
 823
 824 #pragma acc kernels async (1)
 825     {
 826         int ii;
 827
 828         for (ii = 0; ii < N; ii++)
 829             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 830     }
 831
 832 #pragma acc update host (a[0:N], b[0:N], c[0:N]) wait (1)
 833
 834     }
 835
 836     for (i = 0; i < N; i++)
 837     {
 838         if (a[i] != 4.0)
 839             abort ();
 840
 841         if (b[i] != 16.0)
 842             abort ();
 843
 844         if (c[i] != 4.0)
 845             abort ();
 846     }
 847
 848
 849     for (i = 0; i < N; i++)
 850     {
 851         a[i] = 5.0;
 852         b[i] = 0.0;
 853         c[i] = 0.0;
 854         d[i] = 0.0;
 855         e[i] = 0.0;
 856     }
 857
 858 #pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
 859     {
 860
 861 #pragma acc kernels async (1)
 862     {
 863         int ii;
 864
 865         for (ii = 0; ii < N; ii++)
 866             b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
 867     }
 868
 869 #pragma acc kernels async (1)
 870     {
 871         int ii;
 872
 873         for (ii = 0; ii < N; ii++)
 874             c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
 875     }
 876
 877 #pragma acc update host (a[0:N], b[0:N], c[0:N]) async (1)
 878
 879 #pragma acc wait (1)
 880
 881     }
 882
 883     for (i = 0; i < N; i++)
 884     {
 885         if (a[i] != 5.0)
 886             abort ();
 887
 888         if (b[i] != 25.0)
 889             abort ();
 890
 891         if (c[i] != 4.0)
 892             abort ();
 893     }
 894
 895     acc_shutdown (acc_device_nvidia);
 896
 897     return 0;
 898 }