1  
       2  #if ACC_DEVICE_TYPE_nvidia
       3  
       4  #pragma acc routine nohost
       5  static 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  
      16  void
      17  delay (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  
      31  void
      32  delay2 (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  }