1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
2 /* { dg-additional-options "-lcuda" } */
3 /* { dg-require-effective-target openacc_cuda } */
12 main (int argc
, char **argv
)
21 unsigned long *a
, *d_a
, dticks
;
28 acc_init (acc_device_nvidia
);
30 devnum
= acc_get_device_num (acc_device_nvidia
);
32 r
= cuDeviceGet (&dev
, devnum
);
33 if (r
!= CUDA_SUCCESS
)
35 fprintf (stderr
, "cuDeviceGet failed: %d\n", r
);
40 cuDeviceGetAttribute (&nprocs
, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT
,
42 if (r
!= CUDA_SUCCESS
)
44 fprintf (stderr
, "cuDeviceGetAttribute failed: %d\n", r
);
48 r
= cuDeviceGetAttribute (&clkrate
, CU_DEVICE_ATTRIBUTE_CLOCK_RATE
, dev
);
49 if (r
!= CUDA_SUCCESS
)
51 fprintf (stderr
, "cuDeviceGetAttribute failed: %d\n", r
);
55 r
= cuModuleLoad (&module
, "subr.ptx");
56 if (r
!= CUDA_SUCCESS
)
58 fprintf (stderr
, "cuModuleLoad failed: %d\n", r
);
62 r
= cuModuleGetFunction (&delay
, module
, "delay");
63 if (r
!= CUDA_SUCCESS
)
65 fprintf (stderr
, "cuModuleGetFunction failed: %d\n", r
);
69 nbytes
= nprocs
* sizeof (unsigned long);
73 dticks
= (unsigned long) (dtime
* clkrate
);
75 a
= (unsigned long *) malloc (nbytes
);
76 d_a
= (unsigned long *) acc_malloc (nbytes
);
78 acc_map_data (a
, d_a
, nbytes
);
80 kargs
[0] = (void *) &d_a
;
81 kargs
[1] = (void *) &dticks
;
83 for (i
= 0; i
< N
; i
++)
85 streams
[i
] = (CUstream
) acc_get_cuda_stream (i
);
86 if (streams
[i
] != NULL
)
89 r
= cuStreamCreate (&streams
[i
], CU_STREAM_DEFAULT
);
90 if (r
!= CUDA_SUCCESS
)
92 fprintf (stderr
, "cuStreamCreate failed: %d\n", r
);
96 if (!acc_set_cuda_stream (i
, streams
[i
]))
100 for (i
= 0; i
< N
; i
++)
102 r
= cuLaunchKernel (delay
, 1, 1, 1, 1, 1, 1, 0, streams
[i
], kargs
, 0);
103 if (r
!= CUDA_SUCCESS
)
105 fprintf (stderr
, "cuLaunchKernel failed: %d\n", r
);
109 if (acc_async_test (i
) != 0)
111 fprintf (stderr
, "asynchronous operation not running\n");
116 sleep ((int) (dtime
/ 1000.0f
) + 1);
118 for (i
= 0; i
< N
; i
++)
120 if (acc_async_test (i
) != 1)
122 fprintf (stderr
, "found asynchronous operation still running\n");
132 acc_shutdown (acc_device_nvidia
);
137 /* { dg-output "" } */