CUDA 5の新機能(3): CPU処理のStream投入

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

この記事はRC版の情報を基にしており、リリース版では変更が有ることが分かっています。リリース版の仕様に合わせて修正していますが、古い情報が残っている可能性が有ります。

CUDAにはStreamという機能が有り、GPU計算のKernelやメモリ転送を、CPUのメインスレッドとは独立のキューで管理することができる。このStreamの機能がCUDA 5で拡張され、CPU処理をStreamに投入できるようになった。

API

cudaError_t cudaStreamAddCallback(cudaStream_t, cudaStreamCallback_t, void *, unsigned int)
第1引数のStreamのキューに、第3引数を引数にした、関数ポインタである第2引数の呼び出し処理を投入する。第4引数は処理のフラグであり、CPU処理のコールバック関数を呼び出し次第ブロッキングせずにすぐにキューの次の指令の処理に移るcudaStreamCallbackNonblocking(第4引数が0のときのデフォルト値)と、コールバック関数の終了を待ってキューの次の指令を処理するcudaStreamCallbackBlockingとの、2種類の定数が定義されている。第4引数は5.0リリース版では常に0が指定され、常にブロッキング動作する。第2引数の型cudaStreamCallback_tは、具体的にはtypedef void (*cudaStreamCallback_t)(cudaStream_t, cudaError_t, void *)と定義されている関数ポインタであり、第1引数はStream、第2引数は直前のStream処理のエラーコード、第3引数はcudaStreamAddCallbackの第3引数に対応する引数である。

リリース版に於ける使用例

#include <cstdio>
#include <unistd.h>

void
callback1(cudaStream_t stream, cudaError_t status, void *userData)
{
  sleep(10);
  printf("callback1: %c\n", *reinterpret_cast<const char *>(userData));
}

void
callback2(cudaStream_t stream, cudaError_t status, void *userData)
{
  printf("callback2: %c\n", *reinterpret_cast<const char *>(userData));
}

int
main()
{
  const size_t N = 1024 * 1024 * 1024;
  char *h_a, *d_a;
  cudaHostAlloc(reinterpret_cast<void **>(&h_a), N, cudaHostAllocMapped);
  cudaHostGetDevicePointer(reinterpret_cast<void **>(&d_a), h_a, 0);
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  cudaMemsetAsync(d_a, 'A', N, stream);
  cudaStreamAddCallback(stream, callback1, h_a, 0);
  cudaMemsetAsync(d_a, 'B', N, stream);
  cudaStreamAddCallback(stream, callback2, h_a, 0);
  cudaStreamSynchronize(stream);
  cudaStreamDestroy(stream);
  cudaFreeHost(h_a);
  return 0;
}
callback1: A
callback2: B

処理のポイントであるが、Streamを使ったcudaMemsetAsyncでまず配列に「A」を代入し、CPU処理であるcallback1を続いてcudaStreamAddCallbackでStreamに投入する。callback1は10秒スリープするが、これはStreamの処理であるので、後続のcudaMemsetAsyncで配列に「B」を代入する処理は、このcallback1の終了を待つ。そのため、callback1は「A」を出力する。次にcallback2がStreamに投入されるが、やはりStreamの実行順に処理が行われ、先のcudaMemsetAsyncを追い越さないため、callback2は「B」を出力する。

RC版当時の使用例(5.0リリース版では動作しない)

#include <cstdio>
#include <unistd.h>

void
callback1(cudaStream_t stream, cudaError_t status, void *userData)
{
  sleep(10);
  printf("callback1: %c\n", *reinterpret_cast<const char *>(userData));
}

void
callback2(cudaStream_t stream, cudaError_t status, void *userData)
{
  printf("callback2: %c\n", *reinterpret_cast<const char *>(userData));
}

int
main()
{
  const size_t N = 1024 * 1024 * 1024;
  char *h_a, *d_a;
  cudaHostAlloc(reinterpret_cast<void **>(&h_a), N, cudaHostAllocMapped);
  cudaHostGetDevicePointer(reinterpret_cast<void **>(&d_a), h_a, 0);
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  cudaMemsetAsync(d_a, 'A', N, stream);
  cudaStreamAddCallback(stream, callback1, h_a, cudaStreamCallbackNonblocking);
  cudaMemsetAsync(d_a, 'B', N, stream);
  cudaStreamAddCallback(stream, callback1, h_a, cudaStreamCallbackBlocking);
  cudaMemsetAsync(d_a, 'C', N, stream);
  cudaStreamAddCallback(stream, callback1, h_a, cudaStreamCallbackNonblocking);
  cudaStreamAddCallback(stream, callback2, h_a, cudaStreamCallbackNonblocking);
  cudaStreamSynchronize(stream);
  sleep(11);
  cudaStreamDestroy(stream);
  cudaFreeHost(h_a);
  return 0;
}
callback1: B
callback1: B
callback1: C
callback2: C

処理のポイントであるが、Streamを使ったcudaMemsetAsyncでまず配列に「A」を代入し、CPU処理であるcallback1を続いてcudaStreamAddCallbackでStreamに投入する。callback1は10秒スリープして、配列の先頭の要素を出力しようとするが、cudaStreamCallbackNonblockingを付けて呼び出しているので、すかさず次のcudaMemsetAsyncで配列に「B」が代入される。次に再びcallback1がStreamに投入されるが、こちらはcudaStreamCallbackBlockingが付いているので、callback1の処理が終わるまでは、次の処理である「C」を代入するcudaMemcpyAsyncには移らない。そのため、最初のcallback1と次のcallback1が順次処理され、共に「B」を出力する。

「C」を代入するcudaMemcpyAsyncの後に、callback1と、10秒スリープの処理が抜かれているcallback2と、共にcudaStreamCallbackNonblocking付きでStreamに投入しているが、CPU処理のコールバック関数同士は、cudaStreamCallbackNonblockingが指定されていても前の処理の完了を待つようである。そのため、後に呼ばれたcallback2callback1の終了を待ち、callback1callback2の順番で出力が行われる。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:CUDA 5の新機能(3): CPU処理のStream投入

個人用ツール