libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c

   1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
   2 /* { dg-additional-options "-lcuda" } */
   3 /* { dg-require-effective-target openacc_cuda } */
   4
   5 #include <openacc.h>
   6 #include <stdlib.h>
   7 #include "cuda.h"
   8
   9 #include <stdio.h>
  10
  11 #define n 128
  12
  13 int
  14 main (void)
  15 {
  16   CUresult r;
  17   CUstream stream1;
  18   int N = n;
  19   int a[n];
  20   int c[n];
  21
  22   acc_init (acc_device_nvidia);
  23
  24   r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
  25   if (r != CUDA_SUCCESS)
  26     {
  27       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
  28       abort ();
  29     }
  30
  31   acc_set_cuda_stream (1, stream1);
  32
  33   for (int i = 0; i < n; i++)
  34     {
  35       a[i] = 3;
  36       c[i] = 0;
  37     }
  38
  39 #pragma acc data copy (a, c) copyin (N)
  40   {
  41 #pragma acc parallel async (1)
  42     ;
  43
  44 #pragma acc parallel async (1) num_gangs (320)
  45     #pragma acc loop gang
  46     for (int ii = 0; ii < N; ii++)
  47       c[ii] = (a[ii] + a[N - ii - 1]);
  48
  49 #pragma acc parallel async (1)
  50     #pragma acc loop seq
  51     for (int ii = 0; ii < n; ii++)
  52       a[ii] = 6;
  53
  54 #pragma acc wait (1)
  55   }
  56
  57   for (int i = 0; i < n; i++)
  58     if (c[i] != 6)
  59       abort ();
  60
  61   return 0;
  62 }