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

   1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
   2 /* { dg-additional-options "-lcuda" } */
   3
   4 #include <stdio.h>
   5 #include <unistd.h>
   6 #include <openacc.h>
   7 #include <cuda.h>
   8
   9 int
  10 main (int argc, char **argv)
  11 {
  12   CUdevice dev;
  13   CUfunction delay;
  14   CUmodule module;
  15   CUresult r;
  16   CUstream stream;
  17   unsigned long *a, *d_a, dticks;
  18   int nbytes;
  19   float dtime;
  20   void *kargs[2];
  21   int clkrate;
  22   int devnum, nprocs;
  23
  24   acc_init (acc_device_nvidia);
  25
  26   devnum = acc_get_device_num (acc_device_nvidia);
  27
  28   r = cuDeviceGet (&dev, devnum);
  29   if (r != CUDA_SUCCESS)
  30     {
  31       fprintf (stderr, "cuDeviceGet failed: %d\n", r);
  32       abort ();
  33     }
  34
  35   r =
  36     cuDeviceGetAttribute (&nprocs, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
  37                           dev);
  38   if (r != CUDA_SUCCESS)
  39     {
  40       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  41       abort ();
  42     }
  43
  44   r = cuDeviceGetAttribute (&clkrate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
  45   if (r != CUDA_SUCCESS)
  46     {
  47       fprintf (stderr, "cuDeviceGetAttribute failed: %d\n", r);
  48       abort ();
  49     }
  50
  51   r = cuModuleLoad (&module, "subr.ptx");
  52   if (r != CUDA_SUCCESS)
  53     {
  54       fprintf (stderr, "cuModuleLoad failed: %d\n", r);
  55       abort ();
  56     }
  57
  58   r = cuModuleGetFunction (&delay, module, "delay");
  59   if (r != CUDA_SUCCESS)
  60     {
  61       fprintf (stderr, "cuModuleGetFunction failed: %d\n", r);
  62       abort ();
  63     }
  64
  65   nbytes = nprocs * sizeof (unsigned long);
  66
  67   dtime = 200.0;
  68
  69   dticks = (unsigned long) (dtime * clkrate);
  70
  71   a = (unsigned long *) malloc (nbytes);
  72   d_a = (unsigned long *) acc_malloc (nbytes);
  73
  74   acc_map_data (a, d_a, nbytes);
  75
  76   kargs[0] = (void *) &d_a;
  77   kargs[1] = (void *) &dticks;
  78
  79   stream = (CUstream) acc_get_cuda_stream (0);
  80   if (stream != NULL)
  81     abort ();
  82
  83   r = cuStreamCreate (&stream, CU_STREAM_DEFAULT);
  84   if (r != CUDA_SUCCESS)
  85     {
  86       fprintf (stderr, "cuStreamCreate failed: %d\n", r);
  87       abort ();
  88     }
  89
  90   if (!acc_set_cuda_stream (0, stream))
  91     abort ();
  92
  93   r = cuLaunchKernel (delay, 1, 1, 1, 1, 1, 1, 0, stream, kargs, 0);
  94   if (r != CUDA_SUCCESS)
  95     {
  96       fprintf (stderr, "cuLaunchKernel failed: %d\n", r);
  97       abort ();
  98     }
  99
 100   if (acc_async_test (0) != 0)
 101     {
 102       fprintf (stderr, "asynchronous operation not running\n");
 103       abort ();
 104     }
 105
 106   /* Test unseen async-argument.  */
 107   if (acc_async_test (1) != 1)
 108     {
 109       fprintf (stderr, "acc_async_test failed on unseen async-argument\n");
 110       abort ();
 111     }
 112
 113   sleep (1);
 114
 115   if (acc_async_test (0) != 1)
 116     {
 117       fprintf (stderr, "found asynchronous operation still running\n");
 118       abort ();
 119     }
 120
 121   acc_unmap_data (a);
 122
 123   free (a);
 124   acc_free (d_a);
 125
 126   acc_shutdown (acc_device_nvidia);
 127
 128   exit (0);
 129 }
 130
 131 /* { dg-output "" } */