1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
2 /* { dg-additional-options "-lcuda" } */
3 /* { dg-require-effective-target openacc_cuda } */
13 main (int argc
, char **argv
)
22 unsigned long *a
, *d_a
, dticks
;
24 float atime
, dtime
, hitime
, lotime
;
29 acc_init (acc_device_nvidia
);
31 devnum
= acc_get_device_num (acc_device_nvidia
);
33 r
= cuDeviceGet (&dev
, devnum
);
34 if (r
!= CUDA_SUCCESS
)
36 fprintf (stderr
, "cuDeviceGet failed: %d\n", r
);
41 cuDeviceGetAttribute (&nprocs
, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT
,
43 if (r
!= CUDA_SUCCESS
)
45 fprintf (stderr
, "cuDeviceGetAttribute failed: %d\n", r
);
49 r
= cuDeviceGetAttribute (&clkrate
, CU_DEVICE_ATTRIBUTE_CLOCK_RATE
, dev
);
50 if (r
!= CUDA_SUCCESS
)
52 fprintf (stderr
, "cuDeviceGetAttribute failed: %d\n", r
);
56 r
= cuModuleLoad (&module
, "subr.ptx");
57 if (r
!= CUDA_SUCCESS
)
59 fprintf (stderr
, "cuModuleLoad failed: %d\n", r
);
63 r
= cuModuleGetFunction (&delay
, module
, "delay");
64 if (r
!= CUDA_SUCCESS
)
66 fprintf (stderr
, "cuModuleGetFunction failed: %d\n", r
);
70 nbytes
= nprocs
* sizeof (unsigned long);
74 dticks
= (unsigned long) (dtime
* clkrate
);
78 a
= (unsigned long *) malloc (nbytes
);
79 d_a
= (unsigned long *) acc_malloc (nbytes
);
81 acc_map_data (a
, d_a
, nbytes
);
83 stream
= (CUstream
) acc_get_cuda_stream (0);
87 r
= cuStreamCreate (&stream
, CU_STREAM_DEFAULT
);
88 if (r
!= CUDA_SUCCESS
)
90 fprintf (stderr
, "cuStreamCreate failed: %d\n", r
);
94 if (!acc_set_cuda_stream (0, stream
))
99 kargs
[0] = (void *) &d_a
;
100 kargs
[1] = (void *) &dticks
;
104 for (i
= 0; i
< N
; i
++)
106 r
= cuLaunchKernel (delay
, 1, 1, 1, 1, 1, 1, 0, stream
, kargs
, 0);
107 if (r
!= CUDA_SUCCESS
)
109 fprintf (stderr
, "cuLaunchKernel failed: %d\n", r
);
116 atime
= stop_timer (0);
119 hitime
+= hitime
* 0.02;
122 lotime
-= lotime
* 0.02;
124 if (atime
> hitime
|| atime
< lotime
)
126 fprintf (stderr
, "actual time < delay time\n");
137 acc_shutdown (acc_device_nvidia
);
142 /* { dg-output "" } */