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

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