スパコンポエムAdC2020 Day2 Nov/2020 #9? Tianhe-2

どーも、山田です。

スーパーコンピュータポエムアドベントカレンダー2020です。
今日からはヌルくいきますよ。

Day2 Tianhe-2

そいつは突然やってきた。
June/2013。理論ピーク性能27.1PFlops, 実効性能17.5PFlopsのTitanが首位を走るランキングに、理論ピーク性能54.9PFlops, 実効性能33.8PFlopsの性能に強化されて現れた、中国のダークホース…
首位の倍の性能を叩き出したのは、中国国防科技大学が誇るTianhe-2。王者IntelのXeonPhiを抱えて、やつはTop500の首位に躍り出て、そして、三年の長きに渡り首位を独走し続けたのであった……

Knights Corner(KNC) 搭載のスパコンの話をしてくれって言われても、まだKNCのスパコンあったっけ…とか思ってたらないんですけど、今もランキングにある(といっていいのかしら…) のが持ってました。中国の誇るTianhe-2。天河2号といったほうが日本では通りがよさそうですが。
もはや懐かしいIntelのXeonPhi第一世代、Knights Corner。そんなんありましたねーって感じの代物です。シリーズ自体が死にましたからね。挙句、どこぞで開発ボードが$150で投げ売りされていたっていうのは記憶に新し…くはないな…なにせ2014年の出来事であるからして

そんなTianhe-2の話をしたいと思います。
オーバービューについては、かのJack Dongarra先生たちのレポートがあるのでこちらをぜひご覧ください。
https://www.dropbox.com/s/ep3pxq4g9d07yh0/tianhe-2-dongarra-report.pdf?dl=0

ちなみにですね。
実はTianhe-2というのはもうなくて、現在はTianhe-2Aというスーパーコンピュータシステムに改良されています。なんと理論ピーク性能は100PFlops。倍になってる。
このTianhe-2Aについては、同じくDongarra先生がレポートを作られているので、そちらをぜひご参照のこと。2AのMatrix2000っていうのも面白いのでAdCで触れたいですね。
https://www.icl.utk.edu/files/publications/2017/icl-utk-970-2017.pdf

Tianhe-2について

紫色の怪しい光がかっこいいスパコンですねえ。
このカラーバー、負荷量に応じて色変わるらしいですよ。スパコンも映えを狙うご時世ですよ。なおこのシステムは結構前にできている模様。

引用元:https://wired.jp/2013/06/10/chinese-supercomputer-destroys-speed-record-and-will-get-much-faster/

さて、スペックですが、

このサイトに載っているのはすでに2Aになっているので、過去の履歴から参照しましょう。

までが2のスペックです。

使っていたCPUはXeon E5-2692v2のデュアルソケット。サイトのほうには2692と書かれていますが、これv2が正しいと思います。IvyBridge-EP世代と書かれているので。
インターコネクトはTH-Express2という謎のインターコネクトを使っています。これは調べても資料が出てこない謎です。

一応ここに軽いペーパーがあるぐらい。
そしてアクセラレータはIntel XeonPhi 31S1P。これはTianhe-2用の特別モデルです。投入直後だったから、歩留まりが上がらなかったのか、本来61コアあるはずのKNCのコアが57コアになっています。
Tianhe-2はこのXeonPhiによって、その巨大な演算性能を得ていたわけなんですね。

理論ピーク性能は54.9PFlops。HPLの実効性能は33.8PFlops。実効効率としては61%なので、HPLの効率としてはあまり高いとはいいがたいです。消費電力は驚きの17.8MW。
当時、二位のTitanが8.2MWだったことを考えると、なるほど電力突っ込めば倍ぐらいは出るのか、という気持ちにはなったりならなかったりします。
ちなみに電力効率としては割と厳しかった京コンピュータが10.5PFlopsで12.6MWだったことを考えると、京コンピュータよりはいいものの、GPUよりは効率が悪い…というのがわかります。

さて、ではこのTianhe-2というスーパーコンピュータで特徴的なXeonPhiについて、少し語っていきましょう。

XeonPhiというアクセラレータについて

XeonPhiというのは、Intelが打ち出したGPUに対する一つの回答でした。…いや、IntelいつもGPUに回答してんな？

XeonPhiの話をするからには、Larrabeeの話は避けては通れないでしょう。
そもそもLarrabeeというのは、Intelが開発していたディスクリートGPUのことであり、アーキテクチャとしての繋がりは全くないですが、いま絶賛Intelが心血を注いでいる(のか？) Xeのご先祖様といえます。
アーキテクチャとかはWikipediaに詳しいです。何でも書いてあんなWikipedia・・・

さて、Larrabeeは「x86でGPUをやろう」…というのがどこまで本当なのかはわかりませんが、CPUと同様の発想で、SIMDレーン長が十分にあれば、CPUと同様の構造であってもグラフィックスも高速に処理できる、というのが根底にあったと考えられます。
2009年当時のGPUとかはVLIW4とかの時代だったので、そういう発想に到るのも無理はないのかという気もしないでもないですが、ともかく、x86をGPUとして使うプロジェクトは、残念ながら頓挫してしまいました。
このLarrabeeでの知見を再利用して、General Purposeなアクセラレータとして開発されたのが、XeonPhiファミリーであり、Tianhe-2に搭載されたのはその第一世代、開発コードKnights Cornerです。

KNCアーキテクチャ

KNCこと第一世代XeonPhiは、PCIeでホストCPUと接続されるアクセラレータです。
ぶっちゃけPCWatchの後藤さんの記事読んだほうがいいです…
https://pc.watch.impress.co.jp/docs/column/kaigai/558738.html

というのもアレなので一応概要ぐらいは書いておくと、KNCは以下のような特徴を持っています

• ホストとはPCIe x16で接続される
• 4スレッド / コア
• コアの構成はインオーダーで非常にシンプル
• コアには32KBのData / InstのL1キャッシュ、512KBのL2キャッシュを内蔵
• キャッシュは4スレッドで共有
• 命令パイプは2つのデュアルイシュー
• パイプ0はVPU(512bit SIMD), x87、整数演算命令を吐く
• パイプ1は整数演算命令を吐く
• コア間はリングバスで接続
• ボード全体のメモリはモデルによって異なり、6〜16GB GDDR5. 帯域も同様に、240〜352GB/s
• クロックあたり512bit loadが2, storeが1動作する
• ボード上でLinuxが動作しており、SSHでログインが可能
• ボード上で動作するLinuxはメモリをRAMディスクとして使用しているため、ファイルシステムとして使うとRAMが減る(！！！！

512bit SIMDというのは、今日でこそメインストリームにもAVX512なんぞが下りてきておりメジャーになっていますが、当時としてはそんな広いSIMDが来るなんて夢みたい！みたいな感じでした。
ただしAVX512という名称はまだなく、Intel Initial Many Core Instructionsと呼ばれていた時代です。なんせAVX2がようやく出たところでしたからね。

ただ、SIMD命令セットとしては非常によくできていて、ほぼすべての命令がマスクを取るという贅沢な作りになっていました。おかげで非常に書きやすいものでした。
どういうことかというと、512bitものWide-SIMDにもなってくると、演算しない要素というのが出てきてしまいます。512bitともなると倍精度浮動小数点ですら8要素、単精度浮動小数点では16要素にもなるわけで、それだけ巨大なSIMDレジスタともなると、端数処理が大変になってくるわけです。
そもそも配列の端とかで本来は読んで/書いてほしくないところに書いてしまうと、その時点でSEGVなわけで、端の処理というのはSIMDの天敵でもありました。わたくしの師匠曰く、人生の半分は端数処理で使ったというそうで、どういう人生を歩んだのか逆に気になるところではあります(大嘘

例えば単精度add命令は

__m512 _mm512_mask_add_ps (__m512 src, __mmask16 k, __m512 a, __m512 b)

のように定義されており、__mmkas16 kのビットが立っているスロットのみが処理されます。
ほぼすべての命令に、上記のようなマスク入力が可能になっており、スロットごとに処理する/しないを選択することができるようになっています。
これはほぼ同様の概念がArm v8のエクステンションであるSVEにも採用されており、順当な仕組みであるとはいえ、Intelが先陣を切って実装したというのは、Intelの偉大さを示すエピソードな気がします(ほんとか？

プログラミングモデルも同様に結構野心的に作られており、以下のような3つのプログラミングモデルが提唱されていました。

1. XeonPhiのみでコードを実行する
2. ホストの特定の部位をXeonPhiにオフロードする
3. XeonPhiとホストを協調動作させる

といっても3は、ホストとXeonPhiで演算を実行するというより、XeonPhiの実行中にデータのやり取りをするという使い方のほうが多かったと思うので、実質2にプラスしたみたいな感じでしょうね。

1のXeonPhiのみでコードを実行する、というのは

• ボード上でLinuxが動作しており、SSHでログインが可能

というののたまものであり、深く考えることなく、XeonPhi用のバイナリをコンパイルしてファイル送ってSSHで実行できる、というものでした。ただし、先述の通りファイルを使うと使えるメモリの量が減るのであんまりトクはしません。あと、基本的にシステムコールが激遅でした。
なので、性能を求めるのなら2のオフロードを主眼に置く必要がありました。

オフロードですが、OpenMP 4.0の仕様を先取りするかのようなoffloadプラグマが実装されており

    // vector add
#pragma offload target(mic) \
    in(A:length(num)) \
    in(B:length(num)) \
    out(C:length(num))
    {
#pragma omp parallel for
        for(size_t i = 0; i < num; ++i) {
            C[i] = A[i] + B[i];
        }
    }

と書くとオフロードできました。target構文でXeonPhiに対してのオフロードを宣言し、in, out構文で入出力を決めていた感じですね。これはなかなか書き味としては斬新で素敵でした。実はマルチデバイスも対応しており、target(mic:0)とか書くとデバイス番号も指定できるというシロモノでした。これはなかなかすごかった。

……が。
野心的で先進的だったのに、KNCがあんまりウケなかったのは、ひとえに遅かったからです。かなしい。
当時、それなりにKNCのコードを書いてましたが、とにかくいろいろなものが遅かった……
512KBのL2キャッシュというのはそれなり以上に大きいものでしたが、いかんせんコヒーレントを取るので、ちょっと書いただけで隣に波及し始めてバスがえらいことになったりするし、そもそも隣のコア見に行くだけでめっちゃ遅いし、それこそ演算器が回りきらないし、という悲しみでした。あと、なんかスレッドの起動がめちゃくちゃ遅く、スレッド起動だけで数百msecとか消費してくれたので、全然関係ないタイミングでウォーミングアップしてスレッドを起こし、スレッドプールにしておく必要があるとかそういうバッドノウハウばっかりたまる石でした。
Tianhe-2に搭載されていた31S1Pはクロックが1.1GHzだったので、その理論ピーク性能は1.0TFlopsと、理論ピーク性能は立派だったんですけどもね…
ただ、最内周をL1キャッシュに確実にヒットさせるようにしてIMCIで512bit SIMDを全力で書いた時の性能は、確かに目を見張るものがありました。当時のXeonと比較しても数倍の性能を叩き出せたことを覚えています。これは面白かった。

とはいえ。
Intel XeonPhiの最大のライバルはNVIDIA K20とかではなく、Intel Xeonであった、というのはよく言われていたことでしたが、結局普通にプログラムを書くと、Xeonのデュアルソケットとトントンぐらいの性能になってしまうことが多発しており、あんまり突き抜けられなくて、これに最適化するモチベーションがわかない、というのは仕方ないことだったと思います。

これの巻き返しを狙うIntelは、次なるXeonPhiこと、KnightsLanding、そしてさらなる発展形のKnightsHillの開発を進めるのでありましたが――

まとめ

• Tianhe-2について述べました。
• Tianhe-2の主たる演算力である第一世代XeonPhi KnightsCornerについて述べました。
• KnightsCornerは楽しい石だったことが伝わればうれしいです(？
• 明日も山田の担当で、今度こそKnightsLandingの話をしたい…