11

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()明らかに、カーネル内などでホスト機能を使用することは許可されていません。これに代わる良い方法は何ですか？

4

3 に答える 3

22

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() を使用して周波数を照会する必要があります。周波数は動的になる可能性があるため、正確なスピン ループを保証することは困難です。

于 2012-06-27T00:56:06.243 に答える
11

を読み取るループでビジー待機できますclock()。

少なくとも 10,000 クロック サイクル待機するには:

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;

注: これはテストされていません。オーバーフローを処理するコードは、@Pedro によってこの回答から借用されました。動作の詳細については、彼の回答と CUDA C Programming Guide 4.2 のセクション B.10 を参照してくださいclock()。コマンドもありclock64()ます。

于 2012-06-27T00:53:57.213 に答える