/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-lcuda" } */
/* { dg-require-effective-target openacc_cuda } */

#include <stdlib.h>
#include <unistd.h>
#include <openacc.h>
#include <sys/time.h>
#include <stdio.h>
#include <cuda.h>

int
main (int argc, char **argv)
{
  const int N = 1024 * 1024;
  int i;
  unsigned char *h;
  void *d;
  float async, sync;
  struct timeval start, stop;
  CUresult r;
  CUstream s;

  acc_init (acc_device_nvidia);

  h = (unsigned char *) malloc (N);

  for (i = 0; i < N; i++)
    {
      h[i] = i;
    }

  d = acc_malloc (N);

  acc_map_data (h, d, N);

  gettimeofday (&start, NULL);

  for (i = 0; i < 100; i++)
    {
#pragma acc update device(h[0:N])
    }

  gettimeofday (&stop, NULL);

  sync = (float) (stop.tv_sec - start.tv_sec);
  sync += (float) ((stop.tv_usec - start.tv_usec) / 1000000.0);

  gettimeofday (&start, NULL);

  r = cuStreamCreate (&s, CU_STREAM_DEFAULT);
  if (r != CUDA_SUCCESS)
	{
	  fprintf (stderr, "cuStreamCreate failed: %d\n", r);
	  abort ();
	}

  if (!acc_set_cuda_stream (0, s))
	  abort ();

  for (i = 0; i < 100; i++)
    {
#pragma acc update device(h[0:N]) async(0)
    }

  acc_wait_all ();

  gettimeofday (&stop, NULL);

  async = (float) (stop.tv_sec - start.tv_sec);
  async += (float) ((stop.tv_usec - start.tv_usec) / 1000000.0);

  if (async > (sync * 1.5))
    abort ();

  acc_unmap_data (h);

  acc_free (d);

  free (h);

  acc_shutdown (acc_device_nvidia);

  return 0;
}

/* { dg-output "" } */