1
2#if ACC_DEVICE_TYPE_nvidia
3
4#pragma acc routine nohost
5static int clock (void)
6{
7  int thetime;
8
9  asm __volatile__ ("mov.u32 %0, %%clock;" : "=r"(thetime));
10
11  return thetime;
12}
13
14#endif
15
16void
17delay (unsigned long *d_o, unsigned long delay)
18{
19  int start, ticks;
20
21  start = clock ();
22
23  ticks = 0;
24
25  while (ticks < delay)
26    ticks = clock () - start;
27
28  return;
29}
30
31void
32delay2 (unsigned long *d_o, unsigned long delay, unsigned long tid)
33{
34  int start, ticks;
35
36  start = clock ();
37
38  ticks = 0;
39
40  while (ticks < delay)
41    ticks = clock () - start;
42
43  d_o[0] = tid;
44
45  return;
46}
47