ストリームを用いたコンカレントカーネルプログラミングと最適化エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014

Size: px

Start display at page:

Download "ストリームを用いたコンカレントカーネルプログラミングと最適化エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014"

しゅんすけたかにし
5 years ago
Views:

1 ストリームを用いたコンカレントカーネルプログラミングと最適化エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014

2 コンカレントな処理の実行システム内部の複数の処理を平行に実行する CPU GPU メモリ転送カーネル実行複数のカーネル間ストリーム GPU 上の処理キューカーネル実行メモリ転送の並列性実行順序

3 DEFAULT STREAM Stream : GPU 上の処理を管理するキュー無指定の場合 Default (NULL) Stream が使用される Memcpy Host Device Kernel 1 Memcpy Device Host Default (NULL) Stream t

4 CPU/GPU のコンカレンシを考慮していない例 cudamemcpy() は同期的に動作する Memcpy CPU からのタスク発行処理開始のレイテンシ同期待ち CPU 上の待ち処理開始のレイテンシ Memcpy Device Host t GPU 上での実行 Memcpy Host Device Kernel 1 Memcpy Device Host t

5 CPU GPU 間のコンカレンシ非同期コピー (cudamemcpyasync()) を使用 CPU からのタスク発行 CPU 上で別の処理を行う Device Synchronize() t 処理開始のレイテンシ GPU 上での実行 MemcpyAsync Host Device Kernel MemcpyAsync Device Host t

6 同期版非同期版の API memcpy memset 系には同期非同期バージョンがある基本的には非同期版を使用同期版 cudamemcpy() cudamemcpy2d() cudamemcpytosymbol() cudamemcpyfromsymbol() cudamemset() 非同期版 cudamemcpyasync() cudamemcpy2dasync() cudamemcpytosymbolasync() cudamemcpyfromsymbolasync() cudamemsetasync() CUDA Runtime API のリファレンスから抜粋

7 ブロックする処理 cudadevicesynchronize() 同期 API 意図せず Block する可能性があるのはメモリ確保解放 cudamalloc() cudafree() など対応 : あらかじめ確保しておく Pageable Memoryを使用したcudaMemcpy() 系 API 対応 : cudahostalloc() を使用して Pinned Memoryをアロケート

8 デモ : NSIGHT でタイムラインを見る Memcpy Host Device Kernel 1 Default (NULL) Stream Memcpy Device Host t

9 DRIVER QUEUE LATENCY ドライバ内部のキューの処理レイテンシ Windows におけるデバイスドライバのモード WDDM Mode : ディスプレイドライバミリ秒を超えるレイテンシが発生しやすい cudastreamquery(null) でデバイスに処理を流し込む TCC Mode : Tesla Compute Cluster レイテンシは数十 μ 秒

10 ストリームを使用したコンカレントプログラミング

11 ユーザが作成できる STREAM Blocking Stream 例 : cudastreamcreate(&stm); cudastreamcreatewithflags(&stm, cudastreamdefault); Non-blocking Stream 例 : cudastreamcreatewithflags(&stm, cudastreamnonblocking); Prioritized Stream ( 今日は説明しません ) 例 : cudastreamcreatewithpriority(&stm, cudastreamdefault, priority);

12 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

13 NON-BLOCKING STREAM Default Stream に対しても非同期順序は保証されない Blocking Stream 1 Kernel 1 Blocking Stream 2 Kernel 2 Blocking Stream 3 Kernel 3 Default Stream Kernel

14 STREAM に対する同期プリミティブ状態確認同期イベント同期 API cudastreamquery() cudastreamsynchronize() cudastreamwaitevent() 説明 Stream 上の処理が完了しているか確認 Stream 上の処理完了を確認同期 Stream 上でEventを待つ

15 今日のお題フィルタ付き動画プレーヤー動かしてみる

16 フィルタ付き動画プレーヤー 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

17 データフロー (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

18 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

19 メモリ転送をコンカレントに実行 ( オーバーラップ ) 実行時間 : 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

20 カーネルも並列実行実行時間 : 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

21 性能比較 # 実装処理時間短縮分メモリ転送時間カーネル実行時間 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 メモリ転送のオーバーラップ分速くなったカーネルのオーバーラップは今回はちょっとだけ効果あり

22 カーネル実行もオーバーラップする別ストリームで実行順序に依存しないデータの依存性がない効果のある事例については機会を改めて (Hyper-Qも扱いたい) オーバーラップしているオーバーラップしていない

23 まとめ 1. CPU GPU 間のコンカレンシ非同期 API の使用 2. カーネルとメモリ転送のコンカレンシ Dual Copy Engine 3. カーネル間のコンカレンシカーネル間のデータ依存性ストリーム Default / Blocking / Non-blocking Stream

名称 : 日本 GPU コンピューティングパートナーシップ (G-DEP) 所在 : 東京都文京区本郷 7 丁目 3 番 1 号東京大学アントレプレナープラザ, 他工場 URL アライアンスパートナーコアテクノロジーパートナー NVIDIA JAPAN ソリュ

GPUDirect の現状整理 multi-gpu に取組むために G-DEP チーフエンジニア河井博紀 (kawai@gdep.jp) 名称 : 日本 GPU コンピューティングパートナーシップ (G-DEP) 所在 : 東京都文京区本郷 7 丁目 3 番 1 号東京大学アントレプレナープラザ, 他工場 URL http://www.gdep.jp アライアンスパートナーコアテクノロジーパートナー

ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014

ストリームを用いたコンカレントカーネルプログラミングと最適化エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014