1 / 26

GPU 間マイグレーション による効率的な並列実行

GPU 間マイグレーション による効率的な並列実行. 鈴木太一郎 額田彰 松岡聡 東京工業大学. GPU によるアプリケーションの高速化. GPU 高い演算性能 高いメモリバンド幅 アプリケーションによって性質が異なる 例１）ステンシル系 → メモリアクセス律速 = コアは暇 例２）分子動力学系 → 演算律速 = メモリコントローラは暇 GPU リソースの効率的な使用 → 複数アプリケーションの同時実行. 実行時間削減. 単独実行. B. A. A. 同時実行. B.

bardia
Download Presentation

GPU 間マイグレーション による効率的な並列実行

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. GPU間マイグレーションによる効率的な並列実行GPU間マイグレーションによる効率的な並列実行 鈴木太一郎額田彰松岡聡 東京工業大学

  2. GPUによるアプリケーションの高速化 GPU • 高い演算性能 • 高いメモリバンド幅 • アプリケーションによって性質が異なる • 例１）ステンシル系 → メモリアクセス律速 = コアは暇 • 例２）分子動力学系 → 演算律速 = メモリコントローラは暇 • GPUリソースの効率的な使用 →複数アプリケーションの同時実行 実行時間削減 単独実行 B A A 同時実行 B

  3. 複数のGPUアプリの同時実行のために解決しなければならない２つの問題複数のGPUアプリの同時実行のために解決しなければならない２つの問題 • デバイスメモリ不足に陥る可能性 • 各々のアプリケーションが自由にデバイスメモリを使用 • 同時実行するアプリケーションの選択 • 異なる性質を持つアプリケーションの同時実行が特に有効 本発表は１の問題を解決する GPGPUの主流であるCUDAを対象 • Hyper-Q

  4. デバイスメモリ不足の回避 • 予め各アプリのメモリ使用量が既知である場合 →実行可能な組み合わせを選択 →適応できる場合が限られる • 全てのアプリに手を加える事が可能である場合 →協調動作してデバイスメモリを順番に使用 →プログラミングコストが高い ユーザーに代わりこれを透過的に実現する Mobile CUDA実行環境を提案

  5. Mobile CUDAの基本的な動作原理 • ホストメモリ容量>>デバイスメモリ容量 • ホスト⇔デバイス転送速度>>ストレージ転送速度(10倍以上) 　→ホストメモリを活用したswap in/out • 退避場所はホストメモリ • デバイスメモリに入れるだけ入れる(移動はアプリ単位) ホストメモリ アプリ４ アプリ３ アプリ４ アプリ１ アプリ２ アプリ３ 容量不足 デバイスメモリ容量

  6. Mobile CUDAのソフトウェア構成 • Manager • デバイスメモリ使用量を統括的に管理 • MOCU Libraryにアプリケーションの中断、退避、実行再開を指示 • MOCU Library • デバイスメモリ使用状況に変化があるときにManagerと通信する • 常に中断、退避できるようにCUDAリソースを監視 CUDA実行環境 Mobile CUDA実行環境 socketを用いた通信 NVIDIA library = CUDA Runtime library (libcudart.soなど)

  7. Managerによるデバイスメモリ管理 ManagerからGPUが割り当てられる コンテキスト分(64[MB])の予約 • 各アプリケーションごとに管理 • 使用済みデバイスメモリ量(used) • 追加で要求しているデバイスメモリ量(requested) • requested値と空き容量を比較し、十分な空き容量があれば処理を続行、そうでないなら処理の中断および退避を命じる 例 __device__ float devmem[100]; 64 デバイス変数分の予約 64 デバイス変数 コンテキストが作成され、デバイス変数分を確保 64 デバイス変数 cudaMalloc()呼び出し メモリ確保の許可をManagerに申請 used requested 64 デバイス変数 cudaMalloc() Managerから許可が降りる デバイスメモリ確保に成功 64 デバイス変数 cudaMalloc()

  8. デバイスメモリ空き容量の判定 • 空き容量はNVIDIA Management Library(NVML)によって取得 • reserved値 • デバイスメモリ予約量 = requested値の総和 アプリケーション1 used1 requested1 アプリケーション2 requested2 used2 NVMLによって取得される空き容量 デバイスメモリ used1 used2 requested1 requested2 他アプリケーション使用可能領域 reserved値 デバイスメモリ容量

  9. Managerによるスケジューリング …実行中のアプリ • アプリ情報をリストで保持 • デバイスメモリ量の比較はリストの先頭から行う • 例 • アプリ1が終了すると、GPU0に2.0[GB]の空き容量が発生 • リストの先頭から中断中アプリのused+requested値と2.0[GB]を比較 • アプリ2が実行を再開する …中断中のアプリ アプリ1 アプリ2 アプリ3 アプリn 実行GPU 0 ✕ ✕ 1 2.0 0.5 1.0 2.5 used[GB] requested[GB] 0.0 1.5 0.5 0.5

  10. 現時点での実装におけるスケジューリングの制限現時点での実装におけるスケジューリングの制限 • 許容アプリケーション数はGPU×4に設定(暫定的) • GPUをアサインするアプリケーション数を制限 • 許容数を超えたらGPUをアサインしない • 実行アプリケーション過多な状況を回避 • main関数呼び出しよりも前に実行中断 • ホストメモリ不足を回避 • なるべく実行アプリ数が少ないGPUをアサインする

  11. MOCU Libraryによるアプリの中断、退避、実行再開 • ホストにデータを退避 • デバイスメモリを解放 • 他のCUDAリソースを破壊 • 待機 • デバイスメモリを再確保 • 退避したデータから復元 • 他のCUDAリソースを再確保 再開後矛盾なく実行を継続できる必要がある(完全再現) チェックポイント・リスタートライブラリ NVCRと類似の手法を用いる

  12. CUDAリソースの再構築 • ポインタ型 → カプセル化で対応可能 • アドレス型 → 関連CUDA APIの再実行(Replay)により再現 • 確保するアドレスに影響をおよぼす関数呼び出しを記録 • 記録した関数を同じ順序で再呼び出しすることで復元 Replayの対象関数一覧

  13. NVCRとの相違 • CUDA 4.0からRuntime libraryの構造が激変 • Driver APIで実装するとコンテキストが残留 →Runtime APIを実装

  14. CUDA Pinnedメモリ • GPUから高速アクセス可能なホスト側のメモリ領域 • デバイスメモリ空間にマップ可能（=デバイスメモリアドレスを持つ） • 作成には2通りの方法がある • Mobile CUDAによる対応 • cudaHostAlloc()はReplayできない(ホスト側での処理) →MOCU Libraryが内部でvalloc()とcudaHostRegister()に分解 • GPU間を移動する可能性 →cudaHostRegisterPortableフラグを付加

  15. 特殊な動作モード • カーネル内でのmalloc()関数やnewキーワードによる動的なデバイスメモリの確保 →MOCU Libraryによる追跡が不可能 　→排他モード(Exclusive Mode)で実行 MOCU Library • アプリケーション起動時にCUDA Fat Binaryを解析 • mallocシンボルを発見したらExclusive Modeで実行 Manager • アサインされたGPUで実行中の他のアプリケーションを全て退避 • Exclusive Modeのアプリケーションが終了するまで他のアプリケーションをアサインしない

  16. 関連研究 • 浜野ら GPUクラスタにおける省電力タスクスケジューリング • CPUアプリケーションとGPUアプリケーション間の競合による性能低下予測モデルの構築およびモデルに基づいたタスクスケジューリング手法 • Takizawa, el al. CheCUDA: A checkpoint/restart tool for CUDA applications • classライブラリによってDriver APIが使用するリソースをoverride • 使用にはアプリケーションのリコンパイルが必要

  17. 関連研究 • 遠藤らHHRT/MC • 並列プログラムをメモリ階層利用可能とするランタイム • プログラムに手を加える必要有り • CUDA Unified Memory from CUDA 6.0 • ホストとデバイス間で自動的にpage in/outする機能 • pros: プログラマの負担なし　 • cons: アプリ数の増加に伴いパフォーマンス低下

  18. 実験環境 • 1ノード • CUDAのバージョンは5.5

  19. 実験に使用したアプリケーション

  20. 同時実行による実行時間増加率 • 2つのアプリケーションの実行時間増加率をa、bとして (1/a)+(1/b)≧1なら有効であると言える • bandwidth同士、devmem同士以外は削減している 同時実行するアプリ 対象アプリ 例：matrixMulはpcieと同時実行すると、実行時間が1.17倍になる

  21. アプリケーションの実行方法 • スケジューラによる起動 • 実行するアプリケーションはランダムに選択 • それぞれのスケジューラで実行するアプリケーションの順番は同一

  22. 実行時間評価 • 100アプリケーションを実行 • GPU数を1から4へ変化させて実行時間を測定 • GPU数1の時に最も実行時間を削減(19.1%)、GPU数4の時は13.3%の実行時間削減 合計実行時間 MOCUスケジューラ シンプルスケジューラ

  23. 消費エネルギー • 100アプリケーション実行時の消費電力遷移図(4GPU)

  24. スケジューリングの制限 複数アプリケーションの同時実行 実行時間の削減 20アプリケーション実行状況 MOCUスケジューラでの実行 Process # Exclusive Modeによる実行 シンプルスケジューラでの実行 Process # GPU0で実行中 GPU1で実行中 GPU２で実行中 GPU３で実行中 バックアップ中 実行時間[sec]

  25. まとめ • CUDAアプリケーションの同時実行 • 全体実行時間の削減 • デバイスメモリ不足の危険性 • Mobile CUDAの提案 • 透過的な実行環境 • アプリケーションを中断、退避、実行再開することでデバイスメモリ不足を回避 • 4GPU環境で100アプリケーション実行時、実行時間を13.3%、消費エネルギーを2.2%削減

  26. 今後の課題 • TextureやArrayなどへの対応(やるだけ) • アプリケーションの選択 • 性質の異なるアプリケーションを同時実行することが重要 • アプリの情報 • 商用アプリケーションなどある程度既知なもの • ユーザーが情報提供可能なもの • 情報が全く無いもの

More Related