| |
| #if ACC_DEVICE_TYPE_nvidia |
| |
| #pragma acc routine nohost |
| static int clock (void) |
| { |
| int thetime; |
| |
| asm __volatile__ ("mov.u32 %0, %%clock;" : "=r"(thetime)); |
| |
| return thetime; |
| } |
| |
| #endif |
| |
| void |
| delay (unsigned long *d_o, unsigned long delay) |
| { |
| int start, ticks; |
| |
| start = clock (); |
| |
| ticks = 0; |
| |
| while (ticks < delay) |
| ticks = clock () - start; |
| |
| return; |
| } |
| |
| void |
| delay2 (unsigned long *d_o, unsigned long delay, unsigned long tid) |
| { |
| int start, ticks; |
| |
| start = clock (); |
| |
| ticks = 0; |
| |
| while (ticks < delay) |
| ticks = clock () - start; |
| |
| d_o[0] = tid; |
| |
| return; |
| } |