Równoważnik usleep() w jądrze CUDA?

Question

Chciałbym wywołać coś takiego jak usleep() wewnątrz jądra CUDA. Podstawowym celem jest sprawienie, aby wszystkie rdzenie GPU były uśpione lub zajęte przez kilka mililitrów - jest to część niektórych testów poprawności, które chcę wykonać dla aplikacji CUDA. Moja próba robi to poniżej:

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

otrzymuję następujący błąd, gdy próbuję skompilować ten korzystając nvcc:

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

Oczywiste jest, że nie wolno mi użyć funkcji gospodarza, takich jak usleep() wewnątrz jądra. Co byłoby dobrą alternatywą do tego?

Answer 1

Możesz zająć się czekaniem z pętlą, która brzmi: clock().

poczekać co najmniej 10.000 cykli zegarowych:

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; 

Uwaga: To jest niesprawdzone. Kod, który obsługuje przepełnienia został pożyczony od this answer przez @Pedro. Aby uzyskać szczegółowe informacje na temat działania clock(), zapoznaj się z jego odpowiedzią i sekcją B.10 Instrukcji programowania CUDA C 4.2. Istnieje również polecenie clock64().

Answer 2

Można obracać na zegarze() lub clock64(). CUDA SDK concurrentKernels sample robi to co następuje:

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

Polecam używając clock64(). clock() i clock64() są w cyklach, więc będziesz musiał zapytać o częstotliwość za pomocą funkcji cudaDeviceProperties(). Częstotliwość może być dynamiczna, więc trudno będzie zagwarantować dokładną pętlę spinową.