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