libgomp/testsuite/libgomp.oacc-c-c++-common/lib-81.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 *streams, stream;
  21   unsigned long *a, *d_a, dticks;
  22   int nbytes;
  23   float atime, dtime;
  24   void *kargs[2];
  25   int clkrate;
  26   int devnum, nprocs;
  27
  28   acc_init (acc_device_nvidia);
  29
  30   devnum = acc_get_device_num (acc_device_nvidia);
  31
  32   r = cuDeviceGet (&dev, devnum);
  33   if (r != CUDA_SUCCESS)
  34     {
  35       fprintf (stderr, "cuDeviceGet failed: %d\n", r);
  36       abort ();
  37     }
  38
  39   r =
  40     cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
  41                           dev);
  42   if (r != CUDA_SUCCESS)
  43     {
  44       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  45       abort ();
  46     }
  47
  48   r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
  49   if (r != CUDA_SUCCESS)
  50     {
  51       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  52       abort ();
  53     }
  54
  55   r = cuModuleLoad (&module, "subr.ptx");
  56   if (r != CUDA_SUCCESS)
  57     {
  58       fprintf (stderr, "cuModuleLoad failed: %d\n", r);
  59       abort ();
  60     }
  61
  62   r = cuModuleGetFunction (&delay, module, "delay");
  63   if (r != CUDA_SUCCESS)
  64     {
  65       fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
  66       abort ();
  67     }
  68
  69   nbytes = nprocs * sizeof (unsigned long);
  70
  71   dtime = 500.0;
  72
  73   dticks = (unsigned long) (dtime * clkrate);
  74
  75   N = nprocs;
  76
  77   a = (unsigned long *) malloc (nbytes);
  78   d_a = (unsigned long *) acc_malloc (nbytes);
  79
  80   acc_map_data (a, d_a, nbytes);
  81
  82   streams = (CUstream *) malloc (N * sizeof (void *));
  83
  84   for (i = 0; i < N; i++)
  85     {
  86       streams[i] = (CUstream) acc_get_cuda_stream (i);
  87       if (streams[i] != NULL)
  88         abort ();
  89
  90       r = cuStreamCreate (&streams[i], CU_STREAM_DEFAULT);
  91       if (r != CUDA_SUCCESS)
  92         {
  93           fprintf (stderr, "cuStreamCreate failed: %d\n", r);
  94           abort ();
  95         }
  96
  97         if (!acc_set_cuda_stream (i, streams[i]))
  98           abort ();
  99     }
 100
 101   init_timers (1);
 102
 103   kargs[0] = (void *) &d_a;
 104   kargs[1] = (void *) &dticks;
 105
 106   stream = (CUstream) acc_get_cuda_stream (N);
 107   if (stream != NULL)
 108     abort ();
 109
 110   r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
 111   if (r != CUDA_SUCCESS)
 112     {
 113       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
 114       abort ();
 115     }
 116
 117   if (!acc_set_cuda_stream (N, stream))
 118     abort ();
 119
 120   start_timer (0);
 121
 122   for (i = 0; i < N; i++)
 123     {
 124       r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
 125       if (r != CUDA_SUCCESS)
 126         {
 127           fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
 128           abort ();
 129         }
 130     }
 131
 132   acc_wait_all_async (N);
 133
 134   for (i = 0; i <= N; i++)
 135     {
 136       if (acc_async_test (i) != 0)
 137         abort ();
 138     }
 139
 140   acc_wait (N);
 141
 142   for (i = 0; i <= N; i++)
 143     {
 144       if (acc_async_test (i) != 1)
 145         abort ();
 146     }
 147
 148   atime = stop_timer (0);
 149
 150   if (atime < dtime)
 151     {
 152       fprintf (stderr, "actual time < delay time\n");
 153       abort ();
 154     }
 155
 156   start_timer (0);
 157
 158   stream = (CUstream) acc_get_cuda_stream (N + 1);
 159   if (stream != NULL)
 160     abort ();
 161
 162   r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
 163   if (r != CUDA_SUCCESS)
 164     {
 165       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
 166       abort ();
 167     }
 168
 169   if (!acc_set_cuda_stream (N + 1, stream))
 170     abort ();
 171
 172   acc_wait_all_async (N + 1);
 173
 174   acc_wait (N + 1);
 175
 176   atime = stop_timer (0);
 177
 178   if (0.10 < atime)
 179     {
 180       fprintf (stderr, "actual time too long\n");
 181       abort ();
 182     }
 183
 184   start_timer (0);
 185
 186   acc_wait_all_async (N);
 187
 188   acc_wait (N);
 189
 190   atime = stop_timer (0);
 191
 192   if (0.10 < atime)
 193     {
 194       fprintf (stderr, "actual time too long\n");
 195       abort ();
 196     }
 197
 198   acc_unmap_data (a);
 199
 200   fini_timers ();
 201
 202   free (streams);
 203   free (a);
 204   acc_free (d_a);
 205
 206   acc_shutdown (acc_device_nvidia);
 207
 208   exit (0);
 209 }
 210
 211 /* { dg-output "" } */