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

   1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
   2 /* { dg-additional-options "-lcuda" } */
   3
   4 #include <stdio.h>
   5 #include <stdlib.h>
   6 #include <unistd.h>
   7 #include <openacc.h>
   8 #include <cuda.h>
   9 #include "timer.h"
  10
  11 int
  12 main (int argc, char **argv)
  13 {
  14   CUdevice dev;
  15   CUfunction delay;
  16   CUmodule module;
  17   CUresult r;
  18   int N;
  19   int i;
  20   CUstream stream;
  21   unsigned long *a, *d_a, dticks;
  22   int nbytes;
  23   float atime, dtime, hitime, lotime;
  24   void *kargs[2];
  25   int clkrate;
  26   int devnum, nprocs;
  27
  28   devnum = 2;
  29
  30   acc_init (acc_device_nvidia);
  31
  32   devnum = acc_get_device_num (acc_device_nvidia);
  33
  34   r = cuDeviceGet (&dev, devnum);
  35   if (r != CUDA_SUCCESS)
  36     {
  37       fprintf (stderr, "cuDeviceGet failed: %d\n", r);
  38       abort ();
  39     }
  40
  41   r =
  42     cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
  43                           dev);
  44   if (r != CUDA_SUCCESS)
  45     {
  46       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  47       abort ();
  48     }
  49
  50   r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
  51   if (r != CUDA_SUCCESS)
  52     {
  53       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  54       abort ();
  55     }
  56
  57   r = cuModuleLoad (&module, "subr.ptx");
  58   if (r != CUDA_SUCCESS)
  59     {
  60       fprintf (stderr, "cuModuleLoad failed: %d\n", r);
  61       abort ();
  62     }
  63
  64   r = cuModuleGetFunction (&delay, module, "delay");
  65   if (r != CUDA_SUCCESS)
  66     {
  67       fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
  68       abort ();
  69     }
  70
  71   nbytes = nprocs * sizeof (unsigned long);
  72
  73   dtime = 200.0;
  74
  75   dticks = (unsigned long) (dtime * clkrate);
  76
  77   N = nprocs;
  78
  79   a = (unsigned long *) malloc (nbytes);
  80   d_a = (unsigned long *) acc_malloc (nbytes);
  81
  82   acc_map_data (a, d_a, nbytes);
  83
  84   r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
  85   if (r != CUDA_SUCCESS)
  86     {
  87       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
  88       abort ();
  89     }
  90
  91   if (!acc_set_cuda_stream (1, stream))
  92     abort ();
  93
  94   stream = (CUstream) acc_get_cuda_stream (0);
  95   if (stream != NULL)
  96     abort ();
  97
  98   r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
  99   if (r != CUDA_SUCCESS)
 100     {
 101       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
 102       abort ();
 103     }
 104
 105   if (!acc_set_cuda_stream (0, stream))
 106     abort ();
 107
 108   init_timers (1);
 109
 110   kargs[0] = (void *) &d_a;
 111   kargs[1] = (void *) &dticks;
 112
 113   start_timer (0);
 114
 115   for (i = 0; i < N; i++)
 116     {
 117       r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
 118       if (r != CUDA_SUCCESS)
 119         {
 120           fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
 121           abort ();
 122         }
 123     }
 124
 125   if (acc_async_test (0) != 0)
 126     abort ();
 127
 128   /* Test unseen async-argument.  */
 129   if (acc_async_test (1) != 1)
 130     abort ();
 131
 132   acc_wait_async (0, 1);
 133
 134   if (acc_async_test (0) != 0)
 135     abort ();
 136
 137   if (acc_async_test (1) != 0)
 138     abort ();
 139
 140   /* Test unseen async-argument.  */
 141   {
 142     if (acc_async_test (2) != 1)
 143       abort ();
 144
 145     acc_wait_async (2, 1);
 146
 147     if (acc_async_test (0) != 0)
 148       abort ();
 149
 150     if (acc_async_test (1) != 0)
 151       abort ();
 152
 153     if (acc_async_test (2) != 1)
 154       abort ();
 155   }
 156
 157   acc_wait (1);
 158
 159   atime = stop_timer (0);
 160
 161   if (acc_async_test (0) != 1)
 162     abort ();
 163
 164   if (acc_async_test (1) != 1)
 165     abort ();
 166
 167   hitime = dtime * N;
 168   hitime += hitime * 0.02;
 169
 170   lotime = dtime * N;
 171   lotime -= lotime * 0.02;
 172
 173   if (atime > hitime || atime < lotime)
 174     {
 175       fprintf (stderr, "actual time < delay time\n");
 176       abort ();
 177     }
 178
 179   acc_unmap_data (a);
 180
 181   fini_timers ();
 182
 183   free (a);
 184   acc_free (d_a);
 185
 186   acc_shutdown (acc_device_nvidia);
 187
 188   exit (0);
 189 }
 190
 191 /* { dg-output "" } */