Equivalente di usleep() nel kernel CUDA?

Question

Mi piacerebbe chiamare qualcosa come usleep() all'interno di un kernel CUDA. L'obiettivo di base è rendere tutti i core GPU in sleep o busywait per un numero di millesimi: fa parte di alcuni controlli di integrità che voglio fare per un'applicazione CUDA. Il mio tentativo di fare questo è al di sotto:

#include <unistd.h> 
#include <stdio.h> 
#include <cuda.h> 
#include <sys/time.h> 

__global__ void gpu_uSleep(useconds_t wait_time_in_ms) 
{ 
    usleep(wait_time_in_ms); 
} 

int main(void) 
{ 
    //input parameters -- arbitrary 
    // TODO: set these exactly for full occupancy 
    int m = 16; 
    int n = 16; 
    int block1D = 16; 
    dim3 block(block1D, block1D); 
    dim3 grid(m/block1D, n/block1D); 

    useconds_t wait_time_in_ms = 1000; 

    //execute the kernel 
    gpu_uSleep<<< grid, block >>>(wait_time_in_ms); 
    cudaDeviceSynchronize(); 

    return 0; 
} 

ricevo il seguente errore quando provo a compilare questo utilizzando NVCC:

error: calling a host function("usleep") from a __device__/__global__ 
     function("gpu_uSleep") is not allowed 

Chiaramente, io non sono autorizzato a utilizzare una funzione host come usleep() all'interno di un kernel. Quale sarebbe una buona alternativa a questo?

Answer 1

È possibile attendere con un ciclo che legge clock().

di attendere almeno 10.000 cicli di clock:

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; 

Nota: Questa è testato. Il codice che gestisce gli overflow è stato preso in prestito da this answer da @Pedro. Vedere la sua risposta e la sezione B.10 nella Guida alla Programmazione CUDA C 4.2 per dettagli su come funziona clock(). C'è anche un comando clock64().

Answer 2

È possibile ruotare su clock() o clock64(). L'esempio CUDA SDK concurrentKernels effettua quanto segue:

__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; 
} 

Si consiglia di utilizzare clock64(). clock() e clock64() sono in cicli, quindi è necessario interrogare la frequenza utilizzando cudaDeviceProperties(). La frequenza può essere dinamica, quindi sarà difficile garantire un loop spin accurato.