libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c

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