koie blog

- 日時: 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個しかなかったが今はコア毎に用意されている。
• この手法は大きな問題を解くためにメモリが欲くてノード数を増やすことで総メモリ数を稼いでいるユーザにはメリットがある。

16:39

• http://gpu-computing.gsic.titech.ac.jp/node/63
• 主催： GPUコンピューティング研究会
• 日時： 2012年3月5日（月）13:30〜 (雨)
• 場所： 東京工業大学(大岡山) 情報ネットワーク演習室 第2演習室

CUDAはまださわったことがないけど、ディレクティブベースならやってもいいかなとおもって勉強しにいってきた。

会場の演習室はWUXGA(1920x1200)のディスプレイがずらっと並んでいてWindows 7 enterprise。講師の画面のミラーが2人に1台ついてた。豪華だなぁ。飲食禁止。

ひさしぶりに松田さんに合った。HPCをやるのことになりそうなので情報収集に来ているとのこと。

ゴードンベル受賞(2011)研究の紹介 / 青木尊之（東京工業大学・学術国際情報センター）

• ゴードンベル賞というのはスパコンをどれだけ応用したかに対する賞。
• 気象計算では賞をとれないとおもっていた。賞をねらうためにはフェーズフィールド法をつかった。効率がいい。
• 凝固して成長する仮定をシミュレーションするのは2次元がやっとで3次元は1本がやっと(それでも1週間かかってしまったり)
• 数mmのオーダでシミュレーションしないと特性がわからない。
• TSUBANE2.0 4node 1ヶ月 10万円しないらしい Amazonより安いよ。
• GPUが90%の性能をもっている。京は1000億円 vs TSUBAME2.0 32億円。
• フェーズフィールドφ: 0が液体1が固体とか
• 隣接格子なのでメモリ律速のようにみえて、実はキャッシュに乗るのでCPU律速になる。非線形の計算なので計算負荷が重い→性能がでやすい。
• 無理にshared memoryをつかう必要はない。勝手にキャッシュにのる。
• あえて2次元分割
• 今はGPU間の転送ができないので一旦CPUに転送してMPI通信。
• 境界の計算を先にやって通信しつつ、オーバラップして中の計算。。
  • 境界はCPUでやって(GPUで計算すうとメモリコピー必要になる)、中はGPU。
• 64bitから32bitにしたら早くなったのはアドレス計算の関係。
• weak scaling 2PFlopsでた。44.5%の実行性能。

OpenACC の紹介 / 丸山直也（東京工業大学・学術国際情報センター）

• OpenACCのコンパイラを入手できていないので細かな点でまちがいがあるかも
• PGIとかCrayとかのほかに、CAPS:フランスの会社もつくっている。
• PGI Accelerator Compierがベースでよく似ている。対象言語はC/C++,Fortran。
• OpenMPにとりこまれることを目的にしていたが、なかなか大変なので先にOpenACCとしてリリース。
• Crayのコンパイラをつくっている人が... (ききもらした;だれかおしえて)
• ポインタがあると並列化がむつかしい。
• 書き方は ホストプログラム + (GPU)カーネル関数
• OpenACCではCPUとGPUのメモリは分離されているように見せている。(共有メモリモデルではない)
• メモリコピーは自動でやってくれる。無駄なコピーを省くのはチューニング項目。
• CUDAだとすべて明示的に書かないといけない。
• OpenACC用語: SIMD(thread?;CPUのSIMD演算に対応) < worker(thread?;CPUのcoreに対応) < gang(thread block;CPUのsocketに対応)
  • (ちがってるかも)
• #pragma acc cache(list)でshared memoryに置く指令(ヒント)
• NVIDIAもOpenACCに注力しているよう。
• windowsではどう？

HMPPの概要説明 / 小野寺高之（株式会社JCC ギミック）

• CAPS: フランスの会社。7人の研究者。10年の歴史がある。
• HMPP: 最低2行のディレクティブでGPUで実行できるようになる。(性能は別にして)
• IntelのMICにも対応予定。
• source-to-sourceコンパイラ： CPUコードのコンパイラそのままつかえる。
  • hmpp gcc myProgram.c
• pure functionの呼出をGPUにオフロードできる。基本は関数単位で。
• #pragma hmpp LABEL codelet で関数定義
• #pragma hmpp LABEL callsiteで関数呼出
• codelet,target=CでCのコードを生成できる。
• GPUが利用できない場合にはCPUにフォールバックするが環境変数HMPPRT_NO_FALLBACK=1を定義すると止まるようになる。
• HMPPRT_LOG_LEVEL=INFOでメモリアロケートとかも表示する。(デバッグ用)
• aquire/release: ハードウェアの確保/解放
• allocate/free: コードレットのメモリ割当/解放
• advancedload/delegatedstore: データ転送
• コードレットグループを指定すると同じハードウェアを使うようにできる。
• MAPを指定すると同じ変数であるとヒントを与えられる。
• リファレンスマニュアル: 英語のみ
• チュートリアル: スライドなど(2.xは日本語あり、3.xは日本語化中)。スライド以上のものは、いまのところない。
• プログラミングマニュアル:
• OepnACCのサポート: 3月中に出るとアナウンスしているが。
• OpenACCは共通サブセットになっていてPGIのモデルが基本になっていて、他のベンダは拡張機能として入れる。

HMPP によるプログラムの高速化実習

• TSUBAME2にログインして用意されたソースコードにpragma directiveを入れているだけの簡単なおしごと..
• 受付でログインアカウントが書いてある紙をもらっているので
• windowsにログインして
• cygwinを起動して
• sshでTSUBAME2にログインして
• 環境を設定するスクリプトを実行して
• エディタでpragmaを入れていくのだが、ちゃんとソースコード上に「ここにかいてね」マークがあるので間違うことはないが
• 入れるディレクティブが長〜いので打ち間違えリスクが高い。
• 細かくチューニングするために変数を列挙しないといけないのが面倒臭い。
• でもって実行しようとするとエラーになって動かない。どうやら皆が一斉に実行するのでGPUのキャパを越えてしまうらしい。
• hmpp --codelet-requiredは便利なオプション。GPUコードが生成できなかったらコンパイルエラーにする。
• あまりに単調な作業だったので途中で抜けた。(だって頭が働いてないんだもん)