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 /* { 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 "" } */