Transcript GPU********* ***********

GPU間マイグレーション
による効率的な並列実行
鈴木太一郎 額田彰 松岡聡
東京工業大学
1
GPUによるアプリケーションの高速化
GPU
• 高い演算性能
• 高いメモリバンド幅
• アプリケーションによって性質が異なる
• 例１）ステンシル系 → メモリアクセス律速 = コアは暇
• 例２）分子動力学系 → 演算律速 = メモリコントローラは暇
• GPUリソースの効率的な使用
→複数アプリケーションの同時実行
単独実行
同時実行
実行時間削減
A
B
A
B
2
複数のGPUアプリの同時実行のために
解決しなければならない２つの問題
1. デバイスメモリ不足に陥る可能性
• 各々のアプリケーションが自由にデバイスメモリを使用
2. 同時実行するアプリケーションの選択
• 異なる性質を持つアプリケーションの同時実行が特に有効
本発表は１の問題を解決する
GPGPUの主流であるCUDAを対象
• Hyper-Q
3
デバイスメモリ不足の回避
• 予め各アプリのメモリ使用量が既知である場合
→実行可能な組み合わせを選択
→適応できる場合が限られる
• 全てのアプリに手を加える事が可能である場合
→協調動作してデバイスメモリを順番に使用
→プログラミングコストが高い
ユーザーに代わりこれを透過的に実現する
Mobile CUDA実行環境を提案
4
Mobile CUDAの基本的な動作原理
• ホストメモリ容量>>デバイスメモリ容量
• ホスト⇔デバイス転送速度>>ストレージ転送速度(10倍以
上)
→ホストメモリを活用したswap in/out
• 退避場所はホストメモリ
• デバイスメモリに入れるだけ入れる(移動はアプリ単位)
ホストメモリ
アプリ４
アプリ１
アプリ３
アプリ２
デバイスメモリ容量
アプリ４
アプリ３
容量不足
5
Mobile CUDAのソフトウェア構成
• Manager
• デバイスメモリ使用量を統括的に管理
• MOCU Libraryにアプリケーションの中断、退避、実行再開を指示
• MOCU Library
• デバイスメモリ使用状況に変化があるときにManagerと通信する
• 常に中断、退避できるようにCUDAリソースを監視
CUDA実行環境
NVIDIA library = CUDA Runtime library
(libcudart.soなど)
Mobile CUDA実行環境
socketを用いた通信
6
Managerによるデバイスメモリ管理
ManagerからGPUが割り当てられる
コンテキスト分(64[MB])の予約
• 各アプリケーションごとに管理
• 使用済みデバイスメモリ量(used)
• 追加で要求しているデバイスメモリ量
(requested)
• requested値と空き容量を比較し、十分な
空き容量があれば処理を続行、そうでな
いなら処理の中断および退避を命じる
64
デバイス変数分の予約
64
デバイス変数
コンテキストが作成され、デバイス変数分を確保
64
used
例 __device__ float devmem[100];
デバイス変数
cudaMalloc()呼び出し
メモリ確保の許可をManagerに申請
requested
64
デバイス変数
cudaMalloc()
Managerから許可が降りる
デバイスメモリ確保に成功
64
デバイス変数
cudaMalloc()
7
デバイスメモリ空き容量の判定
• 空き容量はNVIDIA Management Library(NVML)によって取得
• reserved値
• デバイスメモリ予約量 = requested値の総和
アプリケーション1
アプリケーション2
デバイスメモリ
used1
used2
used1
requested1
requested2
used2
NVMLによって取得される空き容量
requested1
requested2
他アプリケーション使用可能領域
reserved値
デバイスメモリ容量
8
Managerによるスケジューリング
…実行中のアプリ
…中断中のアプリ
• アプリ情報をリストで保持
• デバイスメモリ量の比較はリストの先頭から行う
•例
アプリ1
アプリ2
アプリ3
アプリn
実行GPU
0
✕
✕
1
used[GB]
2.0
0.5
1.0
2.5
requested[GB]
0.0
1.5
0.5
0.5
1.
2.
3.
アプリ1が終了すると、GPU0に2.0[GB]の空き容量が発生
リストの先頭から中断中アプリのused+requested値と2.0[GB]を比較
アプリ2が実行を再開する
9
現時点での実装における
スケジューリングの制限
• 許容アプリケーション数はGPU×4に設定(暫定的)
• GPUをアサインするアプリケーション数を制限
• 許容数を超えたらGPUをアサインしない
• 実行アプリケーション過多な状況を回避
• main関数呼び出しよりも前に実行中断
• ホストメモリ不足を回避
• なるべく実行アプリ数が少ないGPUをアサインする
10
MOCU Libraryによる
アプリの中断、退避、実行再開
1. ホストにデータを退避
2. デバイスメモリを解放
3. 他のCUDAリソースを破壊
4. 待機
5. デバイスメモリを再確保
6. 退避したデータから復元
7. 他のCUDAリソースを再確保
再開後矛盾なく実行を継続できる必要がある(完全再現)
チェックポイント・リスタートライブラリ NVCRと類似の手法を用いる
11
CUDAリソースの再構築
• ポインタ型 → カプセル化で対応可能
• アドレス型 → 関連CUDA APIの再実行(Replay)により再現
• 確保するアドレスに影響をおよぼす関数呼び出しを記録
• 記録した関数を同じ順序で再呼び出しすることで復元
Replayの対象関数一覧
関数名
説明
cudaMalloc
デバイスメモリを確保
cudaFree
デバイスメモリを解放
cudaHostRegister
ホストメモリをデバイスにpin down
cudaHostUnregister
pin downしたホストメモリを解除
cudaMallocPitch
2次元配列としてデバイスメモリを確保
cudaMalloc3D
3次元配列としてデバイスメモリを確保
12
NVCRとの相違
Mobile CUDA
NVCR
対応するCUDA API
CUDA Runtime
CUDA Runtime & Driver
CUDA Pinnedメモリ対応
○
✕
対応CUDAバージョン
5.0, 5.5
2.3, 3.0
• CUDA 4.0からRuntime libraryの構造が激変
• Driver APIで実装するとコンテキストが残留
→Runtime APIを実装
13
CUDA Pinnedメモリ
• GPUから高速アクセス可能なホスト側のメモリ領域
• デバイスメモリ空間にマップ可能（=デバイスメモリアドレスを持つ）
• 作成には2通りの方法がある
関数名
マップ時付加するフラグ
説明
cudaHostAlloc()
cudaHostAllocMapped
領域の確保とpin downを行う
cudaHostRegister()
cudaHostRegisterMapped
確保済みの領域をpin downする
• Mobile CUDAによる対応
• cudaHostAlloc()はReplayできない(ホスト側での処理)
→MOCU Libraryが内部でvalloc()とcudaHostRegister()に分解
• GPU間を移動する可能性
→cudaHostRegisterPortableフラグを付加
14
特殊な動作モード
• カーネル内でのmalloc()関数やnewキーワードによる動的なデバ
イスメモリの確保
→MOCU Libraryによる追跡が不可能
→排他モード(Exclusive Mode)で実行
MOCU Library
• アプリケーション起動時にCUDA Fat Binaryを解析
• mallocシンボルを発見したらExclusive Modeで実行
Manager
• アサインされたGPUで実行中の他のアプリケーションを全て退避
• Exclusive Modeのアプリケーションが終了するまで他のアプリ
ケーションをアサインしない
15
関連研究
• 浜野ら GPUクラスタにおける省電力タスクスケ
ジューリング
• CPUアプリケーションとGPUアプリケーション間の競合に
よる性能低下予測モデルの構築およびモデルに基づ
いたタスクスケジューリング手法
• Takizawa, el al. CheCUDA: A checkpoint/restart tool
for CUDA applications
• classライブラリによってDriver APIが使用するリソースを
override
• 使用にはアプリケーションのリコンパイルが必要
16
関連研究
• 遠藤ら HHRT/MC
• 並列プログラムをメモリ階層利用可能とするランタイム
• プログラムに手を加える必要有り
• CUDA Unified Memory from CUDA 6.0
• ホストとデバイス間で自動的にpage in/outする機能
• pros: プログラマの負担なし
• cons: アプリ数の増加に伴いパフォーマンス低下
17
実験環境
• 1ノード
• CUDAのバージョンは5.5
CPU
GPU
搭載数
2ソケット
4
種類
Xeon E5-2687W
Tesla K20c
周波数
3.10[GHz]
0.71[GHz]
コア数
8
2496
メモリ
128[GB]
4800[MB]
18
実験に使用したアプリケーション
アプリケーション
実行
時間
[sec]
サイズ[GB]
Mobile CUDAの
オーバーヘッド[%]
説明
matrixMul
(Medium,Small,Large)
73.3,
59.4,
56.8
2.0,
1.5,
2.5
0.37,
0.41,
0.06
CUDA SDKに内包．計算律速．
pcie
(Medium, Small, Large)
54.3,
72.0,
50.7
2.0,
1.5,
2.4
0.49,
0.11,
0.18
2本のベクトル加算．PCI-e通信律速．
bandwidth
(Medium, Small, Large)
45.7,
47.0,
52.8
2.0,
1.5,
2.4,
0.00,
0.42,
0.23
2本のベクトル加算．GPUメモリバンド幅律速．
malloc_in_kernel
30.7
2.4
0.06
2本のベクトル加算．Exclusive Modeで実行．
devmem
61.0
2.0
0.07
2本のベクトル加算．デバイス変数を使用．
map
46.4
2.0
0.11
2本のベクトル加算．pinneメモリをGPUへマップ．
19
同時実行による実行時間増加率
• 2つのアプリケーションの実行時間増加率をa、bとして
(1/a)+(1/b)≧1なら有効であると言える
• bandwidth同士、devmem同士以外は削減している
同時実行するアプリ
対
象
ア
プ
リ
matirxMul
pcie
bandwidth
devmem
map
matrixMul
1.94
1.17
1.24
1.22
5.01
pcie
1.63
1.88
1.04
1.03
2.57
bandwidth
3.97
1.62
2.03
1.98
3.00
devmem
3.78
1.54
1.96
2.07
10.72
map
1.10
1.35
1.33
1.02
1.60
例：matrixMulはpcieと同時実行すると、実行時間が1.17倍になる
20
アプリケーションの実行方法
• スケジューラによる起動
スケジューラ
MOCU Library
説明
MOCUスケジューラ
使用
全てのアプリケーションを同時に起動
シンプルスケジューラ
不使用
常に1GPUで1アプリケーションが実行さ
れているようにアプリケーションを起動
• 実行するアプリケーションはランダムに選択
• それぞれのスケジューラで実行するアプリケーションの
順番は同一
21
実行時間評価
• 100アプリケーションを実行
• GPU数を1から4へ変化させて実行時間を測定
• GPU数1の時に最も実行時間を削減(19.1%)、GPU
数4の時は13.3%の実行時間削減
合計実行時間
シンプルスケジューラ
MOCUスケジューラ
22
消費エネルギー
• 100アプリケーション実行時の消費電力遷移図(4GPU)
MOCUスケジューラ
シンプルスケジューラ
消費電力平均[W]
705.2
622.9
消費エネルギー[MJ]
871
891
23
スケジューリングの制限
複数アプリケーションの同時実行
実行時間の削減
20アプリケーション実行状況
Process #
MOCUスケジューラでの実行
Exclusive Modeによる実行
Process #
シンプルスケジューラでの実行
GPU0で実行中
GPU1で実行中
GPU２で実行中
GPU３で実行中
バックアップ中
実行時間[sec]
24
まとめ
• CUDAアプリケーションの同時実行
• 全体実行時間の削減
• デバイスメモリ不足の危険性
• Mobile CUDAの提案
• 透過的な実行環境
• アプリケーションを中断、退避、実行再開することでデ
バイスメモリ不足を回避
• 4GPU環境で100アプリケーション実行時、実行時間を
13.3%、消費エネルギーを2.2%削減
25
今後の課題
• TextureやArrayなどへの対応(やるだけ)
• アプリケーションの選択
• 性質の異なるアプリケーションを同時実行することが重要
• アプリの情報
• 商用アプリケーションなどある程度既知なもの
• ユーザーが情報提供可能なもの
• 情報が全く無いもの
26