1  /* { dg-do run { target openacc_nvidia_accel_selected } } */
       2  /* { dg-additional-options "-lcuda" } */
       3  /* { dg-require-effective-target openacc_cuda } */
       4  
       5  #include <stdlib.h>
       6  #include <unistd.h>
       7  #include <openacc.h>
       8  #include <sys/time.h>
       9  #include <stdio.h>
      10  #include <cuda.h>
      11  
      12  int
      13  main (int argc, char **argv)
      14  {
      15    const int N = 1024 * 1024;
      16    int i;
      17    unsigned char *h;
      18    void *d;
      19    float async, sync;
      20    struct timeval start, stop;
      21    CUresult r;
      22    CUstream s;
      23  
      24    acc_init (acc_device_nvidia);
      25  
      26    h = (unsigned char *) malloc (N);
      27  
      28    for (i = 0; i < N; i++)
      29      {
      30        h[i] = i;
      31      }
      32  
      33    d = acc_malloc (N);
      34  
      35    acc_map_data (h, d, N);
      36  
      37    gettimeofday (&start, NULL);
      38  
      39    for (i = 0; i < 100; i++)
      40      {
      41  #pragma acc update device(h[0:N])
      42      }
      43  
      44    gettimeofday (&stop, NULL);
      45  
      46    sync = (float) (stop.tv_sec - start.tv_sec);
      47    sync += (float) ((stop.tv_usec - start.tv_usec) / 1000000.0);
      48  
      49    gettimeofday (&start, NULL);
      50  
      51    r = cuStreamCreate (&s, CU_STREAM_DEFAULT);
      52    if (r != CUDA_SUCCESS)
      53  	{
      54  	  fprintf (stderr, "cuStreamCreate failed: %d\n", r);
      55  	  abort ();
      56  	}
      57  
      58    if (!acc_set_cuda_stream (0, s))
      59  	  abort ();
      60  
      61    for (i = 0; i < 100; i++)
      62      {
      63  #pragma acc update device(h[0:N]) async(0)
      64      }
      65  
      66    acc_wait_all ();
      67  
      68    gettimeofday (&stop, NULL);
      69  
      70    async = (float) (stop.tv_sec - start.tv_sec);
      71    async += (float) ((stop.tv_usec - start.tv_usec) / 1000000.0);
      72  
      73    if (async > (sync * 1.5))
      74      abort ();
      75  
      76    acc_unmap_data (h);
      77  
      78    acc_free (d);
      79  
      80    free (h);
      81  
      82    acc_shutdown (acc_device_nvidia);
      83  
      84    return 0;
      85  }
      86  
      87  /* { dg-output "" } */