libgomp/testsuite/libgomp.oacc-c-c++-common/lib-70.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 <stdio.h>
   6 #include <stdlib.h>
   7 #include <unistd.h>
   8 #include <openacc.h>
   9 #include <cuda.h>
  10
  11 int
  12 main (int argc, char **argv)
  13 {
  14   CUdevice dev;
  15   CUfunction delay;
  16   CUmodule module;
  17   CUresult r;
  18   const int N = 10;
  19   int i;
  20   CUstream streams[N];
  21   unsigned long *a, *d_a, dticks;
  22   int nbytes;
  23   float 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 = 200.0;
  72
  73   dticks = (unsigned long) (dtime * clkrate);
  74
  75   a = (unsigned long *) malloc (nbytes);
  76   d_a = (unsigned long *) acc_malloc (nbytes);
  77
  78   acc_map_data (a, d_a, nbytes);
  79
  80   kargs[0] = (void *) &d_a;
  81   kargs[1] = (void *) &dticks;
  82
  83   for (i = 0; i < N; i++)
  84     {
  85       streams[i] = (CUstream) acc_get_cuda_stream (i);
  86       if (streams[i] != NULL)
  87         abort ();
  88
  89       r = cuStreamCreate (&streams[i], CU_STREAM_DEFAULT);
  90       if (r != CUDA_SUCCESS)
  91         {
  92           fprintf (stderr, "cuStreamCreate failed: %d\n", r);
  93           abort ();
  94         }
  95
  96         if (!acc_set_cuda_stream (i, streams[i]))
  97           abort ();
  98     }
  99
 100   for (i = 0; i < N; i++)
 101     {
 102       r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, streams[i], kargs, 0);
 103       if (r != CUDA_SUCCESS)
 104         {
 105           fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
 106           abort ();
 107         }
 108
 109       if (acc_async_test (i) != 0)
 110         {
 111           fprintf (stderr, "asynchronous operation not running\n");
 112           abort ();
 113         }
 114     }
 115
 116   sleep ((int) (dtime / 1000.0f) + 1);
 117
 118   for (i = 0; i < N; i++)
 119     {
 120       if (acc_async_test (i) != 1)
 121         {
 122           fprintf (stderr, "found asynchronous operation still running\n");
 123           abort ();
 124         }
 125     }
 126
 127   acc_unmap_data (a);
 128
 129   free (a);
 130   acc_free (d_a);
 131
 132   acc_shutdown (acc_device_nvidia);
 133
 134   exit (0);
 135 }
 136
 137 /* { dg-output "" } */