Tương đương với việc ngủ() trong hạt nhân CUDA?

Question

Tôi muốn gọi một cái gì đó như usleep() bên trong một hạt nhân CUDA. Mục tiêu cơ bản là làm cho tất cả các lõi GPU ngủ hoặc bận rộn với một số millescond - đó là một phần của một số kiểm tra độ chính xác mà tôi muốn làm cho một ứng dụng CUDA. nỗ lực của tôi lúc làm điều này là dưới đây:

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

tôi nhận được lỗi sau khi tôi cố gắng biên dịch này sử dụng NVCC:

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

Rõ ràng, tôi không được phép sử dụng một chức năng máy chủ như usleep() bên trong hạt nhân. Điều gì sẽ là một lựa chọn tốt cho điều này?

Answer 1

Bạn có thể đợi với vòng lặp đọc số clock().

Để chờ ít nhất 10.000 chu kỳ đồng hồ:

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; 

Lưu ý: Đây là chưa được kiểm tra. Mã xử lý tràn được mượn từ this answer bởi @Pedro. Xem câu trả lời và phần B.10 của mình trong Hướng dẫn lập trình CUDA 4.2 để biết chi tiết về cách hoạt động của clock(). Ngoài ra còn có một lệnh clock64().

Answer 2

Bạn có thể quay trên đồng hồ() hoặc clock64(). Mẫu CUDA SDK concurrentKernels thực hiện như sau:

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

Tôi khuyên bạn nên sử dụng clock64(). clock() và clock64() đang trong chu kỳ, do đó bạn sẽ phải truy vấn tần số bằng cudaDeviceProperties(). Tần số có thể động vì vậy sẽ rất khó để đảm bảo vòng quay chính xác.