- 日時: 2016-08-31 10:00〜16:40
- 場所: 東工大キャンパスイノベーションセンター2F (田町)
- 主催: JST CREST「ポストペタスケール時代のメモリ階層の深化に対応するソフトウェア技術」研究チーム
- 案内
- 資料
初回2014年につづいて2度目の開催。
館内の自動販売機でジュースを100円で売ってて安い。
10:00
オープニング / 遠藤敏夫 (東工大)
10:06
10:07
インテルXeon Phiプロセッサのアーキテクチャとメモリ階層 / 池井満（Intel)
- インテルのマーケティング
- スケーラブル・システム・フレームワーク
- インテル・クラスタ・レディ
- Xeonは4IPC(inst/clk)だがPhiは2IPC。そのかわり省電力。
- TSX HLE(hardware lock elision): ロックプレフィックスがついていても無視して楽観的排他制御するらしい(競合したらやりなおし)。
- VMMを経由せずにVMに直接割り込みが渡されるようになった(SR-IOV,Direct I/O)。
- リソース・ディレクタ・テクノロジで使えるキャッシュ容量の制御ができるようになった。
- パッケージ内にMCDRAM(積層DRAM)が載っている。
- 帯域は広いが途中にバッファが入る分、レイテンシが1clk程度遅くなる。
- 16GB, 450GB/s。
- パッケージ内にメモリを載せることでメモリ階層が1段ひきあげられたかんじになる。
- MCDRAMのメモリモード: 3種類ある。BIOSで切り替える必要がある。numactl -Hでみられる。
- Cache mode: MCDAMは外付けのDDRメモリのキャッシュとして動作する。ダイレクトマップなので注意。
- Flat mode: アドレス空間にMCDRAMとDDRを並べて配置する。アプリが使い分ける。
- Hybrid mode: MCDRAMの一部をcacheして残りをflatにする混在モード。
- Intel Parallel Studio EXでいろいろ調べられる。
- AVX-512CD: 左辺値にインデックスがコリジョンしたものがあっても再実行してくれる。
- マスクレジスタでAVXレジスタの書き換えたくないところを指定できる。
- クラスタモード: コアとタグメモリとMCDRAMの関係
- all-to-all mode: 論理アドレスのハッシュ値からタグメモリを持っているタイルにききにいってMDRAMをR/W。
- quadrant mode: タイルを4分割することでall-to-allに比べてタイル-タグ-MDRAMの距離を短くする。
- sub NUMA mode: 完全に4分割して4CPUにみせる。
- 姫野ベンチのgosa1が自動でリダクション命令を使うようにできないので、pragmaを入れる必要がある。
- sub NUMAモードにしてMPIを動かしてhbw_malloc()でMDRAMをアロケートするようにすると性能がでる。
11:00
- Q: インテル的には機械学習はPhiでやるつもりか?
- A: トレーニングはPhi、スコアリングはXeon。
- Q: (メモリまで含めて)電力測定はできるか?
- A: おそらくできる。ターボがあるので管理する必要があるため。メモリについては分からない。
- Q: 半精度演算(16bit float)は入るか?
- A: 今はない。
- Q: どういうツールがあるとうれしいか?
- A: 機械学習系ではCUDAが先行しているのでコミュニティ系のものをやれるとうれしい。
11:08
11:09
Pascal世代以降のGPUメモリシステム / 成瀬彰（NVIDIA）
- Tesla P100 (コードネーム GP100)
- HBM2: ピーク720GB/s、実効で7割は出る。4層で16GB積んでいる。スペースは8層分あるので放熱のためスペーサをかましてある。
- GPUはメモリが少ないとの批判がある→Unified Memory
- Unified Memory:
- CPUメモリとGPUメモリでポインタを書き分けなくてもよくなる。
- cudaMallocManaged()したときに物理的にメモリを確保していたのが遅延割り当てになる。
- GPUメモリをCPUが更新するとGPUカーネルで使うかどうかわからないけど全部コピーしていたのがon-demandコピーになる。リンクリストのようなデータ構造で有効。
- ページ単位でコピーするがページサイズは非公開らしい。
- GPUとCPUが競合しないようにページを保護する仕組みは用意してない(簡単なんだからユーザが自分で用意しろというスタンスらしい)。
- Q: readのためにGPU→CPUにコピーしたページもGPUに書き戻される?
- A: …
- GPUメモリがあふれたらpage outできる。
- ページフォルトのレイテンシは10usくらい。同時に1000程度のページフォルトをあつかえるのでレイテンシは隠せる(との主張だがプログラムはブロックされちゃうよね?との質問もあった)。
- out-of-core
- GPUメモリに載りきらないデータをCPUメモリに載せて演算。
- PFNのChainerはよくできている。NumPyのAPIでCuPyというのをつくっている。
- NVLINK
- CPU-GPU間はPCIeでは遅すぎる。
- POEWR8ではNVLINKが載ってくる。40GB/s。 (PCIeは16GB/s)
- それ以外のCPUではGPU-GPU間のみNVLINKということになる。
12:02
- Q: Chainerがunified memoryに対応したら速くなる?
- A: 速くなるのではなくて、いままで動かなかった規模のものが動くようになる。
- Q: 遠藤先生のところのswap in/outがいらなくなる?
- A: 動きのわかっているアプリならon-demandよりも明示的に記述した方がオーバヘッドが少なくなるので棲み分け。
12:06
昼休み
牛タン圭助へ。混んでて15分くらい待ち。日替わりランチが売り切れたので「牛ハラミ定食」に。野菜少ない。ちょっと味付け濃いのでごはんおかわり。
13:11
New 3D Flash Technologies Offer Both Low Cost and Low Power Solutions / 大島成夫（東芝）
- 2020年の予想: 生成されるデータが44[ZB]なのに対して、供給されるストレージ容量は6.6[ZB]しかない。
- SSDはパフォーマンス重視・容量重視・価格重視に分かれていくと予想される。
- メモリ階層のアクセスタイムを光速で距離に換算して、CPUをシェフに例えるとわかりやすい。
- BiCS(ビックス)
- 三次元スタック構造。
- 従来の平面のNAMDを縦に立てて横に並べた感じになる(ビットライン・ワードライン的には)。
- 2017年内には1Tbit品が出る予定。
- ワイヤ・ボンディングからTSV(through-silicon via)にすることでワイヤーのLCRがなくなるので、信号のアイパターンが開いて1.2Gbpsでもクリア。また消費電力はチップレベルで1/2に、コンポーネントレベルで30%減。
- 2D→BiCSにすることで溜め込める電子数が6倍になるので、2Dでは8値(TLC)が限界だったが16値(QLC)までいけそう(試作あり)。
- SSDはHDDの2〜3倍程度の価格になればデータセンターなどで置き換えが一気に進むと思われる。
13:52
- Q: コントローラの話がなかったが?
- A: 重要。電力やスピード(レイテンシ)の問題(トレードオフ)があるのでむつかしい。誤り訂正をがんばると遅くなるなど。IPを開発中。
- Q: TSVについて。
- A: DRAMのTSVは高価だが、BiCSのTSVは低コスト。
- Q: どこで儲けていく?
- A: intelは3D XpointでCPU寄りをねらっている。東芝は容量寄りで。
- Q: ストレージシステムへの戦略は?
- A: 部品供給側よりもPure Storageのような上位の方が儲かるが、1段ずつfoot chainを上がっていくしかない。
- Q: コントローラは進化しているか?
- A: 少しずつかわっている。詳しくはNDAむすんで。
- Q: 生Flashメモリの標準化はしないのか?
- A: NAND flashは標準化しなかったので(各社自由にやれて)伸びた。違いを楽しんでください!
- Q: BiCSで高さを伸ばすとセルは小さくなる?
- A: 積層数を増やしてもセルは小さくならない。配線遅延はそもそもNANDの動作は遅いので問題にならない。
- Q: endurance重視の製品の予定は?
- A: コントローラである程度なんとかなる。実はエンターブライズ向けとコンシュマ向けとでメモリチップはほぼ同じでコントローラが違うだけ。それで寿命ば倍もちがったりする。寿命をのばしたかったかったら、シーケンシャルアクセスして書き込みはページサイズに合わせるべし。データセンターの人たちならやってくれるはずだ。ドライブレベルでならアプリケーションからヒント情報を与えて寿命をのばすことはできるかもしれない。実際、appleやハイエンド向けではファームウェアのチューニングをしている。Host Managed(コントローラをホスト側にもってくる)という流れもあるがSSDメーカとしては付加価値的に歓迎できない。
14:14
(休憩)
14:24
GDDR・DDR・Flashの多階層メモリを利用するランタイムライブラリと大規模ステンシルへの応用 / 遠藤敏夫 (東工大)
- エクサ時代に向けての目標: 100PB/s, 100PB
- 既存のHPCアプリに手を入れて out-of-core 実行を実現した。
- GPUクラスタといえば CUDA + MPI
- ステンシル計算: 流体シミュレーションで重要な計算カーネル。メモリ・インテンシブなのでGPU向けだがメモリが足りないことも。
- HHRTライブラリ: CUDAとMPIのラッパを提供。 https://github.com/toshioendo/hhrt
- cudaMallocをラップしてメモリ割り当てを掌握、MPIが待ちに入るときにメモリのswapを実行することでGPUメモリ以上のサイズの計算を実行可能に。
- ただ、それだけだと性能が1/30に落ちてしまう。
- テンポラル・ブロッキング(時間方向のブロック化): 時間発展の計算を1空間ブロックに対してまとめて実行することで局所性向上を狙う。
- 55%くらいの性能低下におさえられた。
15:04
- Q: 頻繁にswapしたら意味ない。unified memoryとの統合は?
- A: unified memoryだけでは、まだCPUのメインメモリを越える規模の計算はできないので、いちおう棲み分けはできている。統合はできたらしたい。
- Q: テンポラルブロッキングについて
- A: コード1万行のうち500行くらい。時間ステップは50 stepくらいをやっている。(これは計算による)
- BLASでもテンポラルブロッキングはやっている
15:12
15:13
動的バイナリ変換によるメモリ階層性能プロファイリングと透過的メモリ階層チューニング / 佐藤幸紀 (東工大)
- openMPやMPIを使うだけでは性能は出ない。ノード単体のメモリチューニングが必須。チューニングは職人芸。
- 開発:
- Exana: プロファイリング
- ExanaDBT: 透過的チューニング
- Exana:
- ループ構造の階層を表示
- メモリ帯域やアクセス回数を表示
- キャッシュの挙動はシミュレーションに基づく。(100倍程度の速度低下)
- アクセスパターン・ワーキングセットを調べられる
- ExanaDBT:
- バイナリをmcsemaでLLVM IRに変換。
- Pollyでループ最適化。
- Pin tool setでバイナリ書き換え。
15:42
- Q: データ依存解析の出力は表示が不十分なのでは?
- A: 出力をみてパイプライン構造が見えてくるのはチューニング屋さんにはメリットがあるが、アプリを作った人には自明。
- Q: ワーキングセットツールの位置付けは?
- A: ブロッキングサイズを調べるのにつかえる。
- Q: DBTとPGOの比較は?
- A: -O3ならやや効果があるとおもっている。
- Q: HPCではホットスポット最適化の要望が多い?
- A: ソースコード再コンパイルでよいという声も多い。
- Q: x86だけなのはHPC的によくないのでは?
- A: LLVMのコミュニティでPollyが対応していくと期待。
- そもそもPollyはきれいなループ構造でないと最適化してくれないとかループの中に関数呼び出しがあるとダメとか、いろいろヘボい。
- キャッシュのシミュレーションがFIFOなのはよくないのでは?
15:56
15:56
Flash利用によるout-of-coreステンシルアルゴリズムとブロックサイズ自動チューニングシステム / 緑川博子 (成蹊大)
- ioDrive高すぎ
- メモリに載りきらない問題サイズをSSDに逃す
- DRAMとSSDの速度差1000倍くらいある。
- スペーシャルブロッキングだけだと64倍遅くなる。
- テンポラルブロッキングもやると1.3倍におさまる。(実用的)
- ブロックサイズの決定は、問題サイズとデバイス情報を与えることで計算により算出する。いくつか試行して最適なものを選ぶ、というのではない。
- swapやmmapではコピーする時間を隠しきれなかったが、kernel版aioをつかうことでほぼ隠しきれた。内部では細かくaioを出すのではなく、まずブロックの半分をaioでとってきて計算開始、裏で残り半分をaioでとってくる構成におちついた(細かくやっても手間のわりに効果なし)。
16:31
- Q: SSDではなくHDDでも効果があるか?
- A: シーケンシャルアクセスにならない(ベージバウンダリにするためにいろいろ詰め物してる)のでダメだとおもう。SSDとHDDでは性能が1000倍くらい違うし。
- linuxのI/Oサブシステムが昔はキューが1個しかなかったが今はコア毎に用意されている。
- この手法は大きな問題を解くためにメモリが欲くてノード数を増やすことで総メモリ数を稼いでいるユーザにはメリットがある。
