Equivalente de usleep() en kernel CUDA?

Question

Me gustaría llamar a algo como usleep() dentro de un kernel CUDA. El objetivo básico es hacer que todos los núcleos de la GPU duerman o queden ocupados durante varios miles de segundos; es parte de algunas comprobaciones de cordura que quiero hacer para una aplicación CUDA. Mi intento de hacer esto es a continuación:

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

me sale el siguiente error al intentar compilar este usando NVCC:

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

Claramente, no se me permite utilizar una función como anfitrión usleep() dentro de un kernel. ¿Cuál sería una buena alternativa a esto?

Answer 1

Puedes esperar ocupado con un ciclo que dice clock().

que esperar por lo menos 10.000 ciclos de reloj:

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: Esto no se ha probado. El código que maneja los desbordamientos fue tomado prestado de this answer por @Pedro. Consulte su respuesta y la sección B.10 en la Guía de programación de CUDA C 4.2 para obtener detalles sobre cómo funciona clock(). También hay un comando clock64().

Answer 2

Puede activar el reloj() o el reloj64(). El ejemplo de CUTA SDK concurrentKernels hace esto hace lo siguiente:

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

Recomiendo usar clock64(). clock() y clock64() están en ciclos, por lo que tendrá que consultar la frecuencia usando cudaDeviceProperties(). La frecuencia puede ser dinámica, por lo que será difícil garantizar un ciclo de giro preciso.