Equivalent of usleep() in CUDA kernel?

You can busy wait with a loop that reads clock().

To wait for at least 10,000 clock cycles:

clock_t start = clock();
clock_t now;
for (;;) {
  now = clock();
  clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
  if (cycles >= 10000) {
    break;
  }
}
// Stored "now" in global memory here to prevent the
// compiler from optimizing away the entire loop.
*global_now = now;

Note: This is untested. The code that handles overflows was borrowed from this answer by @Pedro. See his answer and section B.10 in the CUDA C Programming Guide 4.2 for details on how clock() works. There is also a clock64() command.

You can spin on clock() or clock64(). The CUDA SDK concurrentKernels sample does this does the following:

__global__ void clock_block(clock_t *d_o, clock_t clock_count)
{
    clock_t start_clock = clock();
    clock_t clock_offset = 0;
    while (clock_offset < clock_count)
    {
        clock_offset = clock() - start_clock;
    }
     d_o[0] = clock_offset;
}

I recommend using clock64(). clock() and clock64() are in cycles so you will have to query the frequency using cudaDeviceProperties(). The frequency can be dynamic so it will be hard to guarantee an accurate spin loop.

With recent versions of CUDA, and a device with Compute Capability 7.0 or later (Volta, Turing, Ampere etc.), you can use the __nanosleep() primitive:

void __nanosleep(unsigned ns);

which obviates the need for busy-sleeping as suggested in older answers.