CUDAカーネルのusleep（）に相当しますか？

Question

私はusleep()のようなものをCUDAカーネルの中で呼びたいと思います。基本的な目標は、すべてのGPUコアを数ミリ秒間スリープ状態またはビジーウェイト状態にすることです。これは、CUDAアプリケーションのためにやってみたいいくつかのサニティチェックの一部です。これを行うに私の試みは以下の通りです：

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

私はNVCCを使用して、これをコンパイルしようとすると、私は次のエラーを取得する：

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

明らかに、私はのようなホスト機能を使用することはできませんよusleep()はカーネル内にあります。これに代わる良い選択肢は何でしょうか？

Answer 1

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; 

注：これはテストされていない少なくとも10,000クロックサイクルを待つために

。オーバーフローを処理するコードは@Pedroによってthis answerから借用されました。 clock()の動作の詳細については、CUDA Cプログラミングガイド4.2の回答とセクションB.10を参照してください。 clock64()コマンドもあります。

Answer 2

clock（）またはclock64（）でスピンすることができます。 CUDA SDKのconcurrentKernelsサンプルでは次のようになります。

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

clock64（）を使用することをお勧めします。 clock（）とclock64（）は周期的なので、cudaDeviceProperties（）を使用して周波数を問い合わせる必要があります。周波数は動的であるため、正確なスピンループを保証することは困難です。