ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014
コンカレントな処理の実行 システム内部の複数の処理を 平行に実行する CPU GPU メモリ転送 カーネル実行 複数のカーネル間 ストリーム GPU 上の処理キュー カーネル実行 メモリ転送の並列性 実行順序
DEFAULT STREAM Stream : GPU 上の処理を管理するキュー 無指定の場合 Default (NULL) Stream が使用される Memcpy Host Device Kernel 1 Memcpy Device Host Default (NULL) Stream t
CPU/GPU のコンカレンシを考慮していない例 cudamemcpy() は 同期的に動作する Memcpy CPU からのタスク発行 処理開始のレイテンシ同期待ち CPU 上の待ち 処理開始のレイテンシ Memcpy Device Host t GPU 上での実行 Memcpy Host Device Kernel 1 Memcpy Device Host t
CPU GPU 間のコンカレンシ 非同期コピー (cudamemcpyasync()) を使用 CPU からのタスク発行 CPU 上で別の処理を行う Device Synchronize() t 処理開始のレイテンシ GPU 上での実行 MemcpyAsync Host Device Kernel MemcpyAsync Device Host t
同期版 非同期版の API memcpy memset 系には 同期 非同期バージョンがある 基本的には 非同期版を使用 同期版 cudamemcpy() cudamemcpy2d() cudamemcpytosymbol() cudamemcpyfromsymbol() cudamemset() 非同期版 cudamemcpyasync() cudamemcpy2dasync() cudamemcpytosymbolasync() cudamemcpyfromsymbolasync() cudamemsetasync() CUDA Runtime API のリファレンスから抜粋
ブロックする処理 cudadevicesynchronize() 同期 API 意図せず Block する可能性があるのは メモリ確保 解放 cudamalloc() cudafree() など 対応 : あらかじめ確保しておく Pageable Memoryを使用したcudaMemcpy() 系 API 対応 : cudahostalloc() を使用して Pinned Memoryをアロケート
デモ : NSIGHT でタイムラインを見る Memcpy Host Device Kernel 1 Default (NULL) Stream Memcpy Device Host t
DRIVER QUEUE LATENCY ドライバ内部のキューの処理レイテンシ Windows におけるデバイスドライバのモード WDDM Mode : ディスプレイドライバ ミリ秒を超えるレイテンシが発生しやすい cudastreamquery(null) で デバイスに処理を流し込む TCC Mode : Tesla Compute Cluster レイテンシは 数十 μ 秒
ストリームを使用したコンカレントプログラミング
ユーザが作成できる STREAM Blocking Stream 例 : cudastreamcreate(&stm); cudastreamcreatewithflags(&stm, cudastreamdefault); Non-blocking Stream 例 : cudastreamcreatewithflags(&stm, cudastreamnonblocking); Prioritized Stream ( 今日は説明しません ) 例 : cudastreamcreatewithpriority(&stm, cudastreamdefault, priority);
BLOCKING STREAM Default Stream と同期する カーネル間でデータの依存性がある場合に有効 順序は保証されない Blocking Stream 1 Kernel 1 Blocking Stream 2 Kernel 2 Blocking Stream 3 Default Stream Kernel 3 Kernel 1-3 を待つ Kernel
NON-BLOCKING STREAM Default Stream に対しても非同期 順序は保証されない Blocking Stream 1 Kernel 1 Blocking Stream 2 Kernel 2 Blocking Stream 3 Kernel 3 Default Stream Kernel
STREAM に対する同期プリミティブ 状態確認 同期 イベント同期 API cudastreamquery() cudastreamsynchronize() cudastreamwaitevent() 説明 Stream 上の処理が完了しているか確認 Stream 上の処理完了を確認 同期 Stream 上でEventを待つ
今日のお題 フィルタ付き動画プレーヤー 動かしてみる
フィルタ付き動画プレーヤー 3 CPU スレッドでパイプラインを構成 Decoder Thread Decode (YV12) Windows Media Foundation Processing Thread YV12 RGB (half float) Convolution Filter R G B Draw to OpenGL Tex CUDA App Main Thread OpenGL Draw to Window
データフロー (PROCESSING THREAD) 入力 (CPU) YV12 Frame GPU YV12 RGB (half float) RGB (3 plane) Half float Convolution Filter R G B RGB (3 plane) Half float Draw to OpenGL Tex 出力 (CPU) Host Buffer R G B
DEFAULT STREAM を使用した場合 実行時間 : 4.17 ms データ転送 : 1.42 ms カーネル : 2.43 ms Memcpy Host Device Convolution Filter Default Stream NV12 YV12 RGB R G B Draw to OpenGL Tex Memcpy Device Host R G B
メモリ転送をコンカレントに実行 ( オーバーラップ ) 実行時間 : 2.85 ms データ転送 : 1.25 ms カーネル : 2.42 ms Non-blocking Stream 1 Non-blocking Stream 2 Memcpy Host Device NV12 Memcpy Device Host R G B Default Stream YV12 RGB Convolution Filter R G B Draw to OpenGL Tex
カーネルも並列実行 実行時間 : 2.78 ms, データ転送 : 1.15 ms, カーネル : 2.37 ms Non-blocking Stream 1 Non-blocking Stream 2 Memcpy Host Device NV12 Memcpy Device Host R G B Convolution Filter Blocking Stream 1 Blocking Stream 2 Blocking Stream 3 Default Stream YV12 RGB R G B Draw to OpenGL Tex
性能比較 # 実装処理時間短縮分メモリ転送時間 カーネル実行時間 1 Default Stream のみ 4.17 ms ー 1.42 ms 2.43 ms 2 メモリ転送をオーバーラップ 3 カーネル実行もオーバーラップ 2.85 ms 1.32 ms 1.25 ms 2.42 ms 2.78 ms 1.39 ms 1.15 ms 2.37 ms メモリ転送のオーバーラップ分 速くなった カーネルのオーバーラップは 今回は ちょっとだけ効果あり
カーネル実行もオーバーラップする 別ストリームで実行 順序に依存しない データの依存性がない 効果のある事例については 機会を改めて (Hyper-Qも扱いたい) オーバーラップしている オーバーラップしていない
まとめ 1. CPU GPU 間のコンカレンシ 非同期 API の使用 2. カーネルとメモリ転送のコンカレンシ Dual Copy Engine 3. カーネル間のコンカレンシ カーネル間のデータ依存性 ストリーム Default / Blocking / Non-blocking Stream