0130_FrontISTR研究会_V3

Transcription

1 Intel Xeon Phi (Knights Landing) のパフォーマンス評価の 例 東京 学 学院 新領域創成科学研究科 松 和, 森 直樹, 奥 洋司 2017 年 1 30 第 33 回 FrontISTR 研究会 2017/1/30 FrontISTR 研究会 1

2 次 背景と 的 KNLのアーキテクチャ メモリモードとクラスタモード STREAM triadによる性能評価 FrontISTRによる性能評価 まとめ 参考 献 2017/1/30 FrontISTR 研究会 2

3 背景と 的 背景 計算機の単体性能向上において, 動作周波数の向上が限界に達する 消費電 の増, 発熱の限界 2004 年ごろから計算機はマルチコア化, 並列処理の流れ 現在のスパコンも並列処理で性能向上を実現 近年は Multi-Channel DRAM(MCDRAM) のような 速メモリが CPU に搭載される 的 XeonPhi Knights Landing ( 以後 KNL) のアーキテクチャ理解 STREAM ベンチマーク, 並列有限要素法ソルバー FrontISTR を いた性能評価 2017/1/30 FrontISTR 研究会 3

4 KNLのアーキテクチャ KNLにはDDR4と積層メモリ (MCDRAM) が搭載 図1-2 KNLのプリント基板 Tile up to 36 Tiles 72 cores 図1-3 Tileの概観[2] 図1-1 KNL Package [1] VPU(vector processing unit) CHA(Caching/Home Agent) DDR4のメモリサイズ MCDRAMのメモリサイズ 512bit register 機能は後で詳しく紹介 64GB 6channelにより最 384GB まで搭載可 2GB 8個で計16GBが搭載 [1] Avinash Sodani, Intel Xeon Phi Processor Knights Landing Architectural Overview, [2] Avinash Sodani, Knights Landing (KNL):2nd Generation Intel Xeon Phi Processor, /1/30 FrontISTR研究会 4

5 メモリモードとクラスタモード メモリモード, クラスタモードという合計 9 通りの使い がある メモリモード (DDR4,MCDRAM の使い を決める ) クラスタモード ( タイル, メモリ間の通信 法を決める ) Flat Cache Hybrid ( ) All to All ( ) Quadrant/Hemisphere SNC(Sub-NUMA Clustering) ( ) 2 つのモードは BIOS で設定し,reboot することで変更できる ( )Oakforest-PACS では選択不可 2017/1/30 FrontISTR 研究会 5

6 メモリモード メモリ (DDR4,MCDRAM) の使い に 3 つのモードがある Flat Cache Hybrid Flat モード Cache モード Hybrid モード MCDRAM は NUMA ノードとして使われる MCDRAM の使 はユーザーが制御 16GB MCDRAM DDR or Physical Address MCDRAM を Last Level Cache(LLC) として使 ハードウェアが MCDRAM の使い を制御 16GB MCDRAM DDR 4or8GB MCDRAM Flat モードと Cache モードの組み合わせ メモリ 率は設定可能 8 or 12GB MCDRAM DDR or Physical Address 図 2-1 Flat モードの概要図 2-2 Cache モードの概要図 2-3 Hybrid モードの概要 メモリモード変更は BIOS で設定し,reboot することで有効になる 2017/1/30 FrontISTR 研究会 6

7 クラスタモード タイル, メモリ間の通信 法に きく 3 つのモードがある All to All Quadrant/Hemisphere SNC(Sub-NUMA Clustering) 図 1-3 Tile の概観 CPU コアは, 欲しいデータがどこのタイル (L2 キャッシュ ) にあるかを Caching/Home Agent(CHA) に確認 L2 miss した時にクラスタモードの性能差がでる 図 3 L2 miss 時のメモリアクセスの流れ Typical Read L2 miss 1. L2 miss encountered 2. Send request to the distributed directory 3. Miss in the directory. Forward to memory 4. Memory send the data to the requestor 2017/1/30 FrontISTR 研究会 7

8 クラスタモード : All to All NUMA(Non-Uniform Memory Access) Tile のディレクトリとメモリに affinity がない Tile のディレクトリ配置とメモリ通信はハードウェアに 動で任せる DDR4 $numactl --membind=0./a.out MCDRAM $numactl --membind=1./a.out 図 4 All to All の概念 3 つのモードのうち 番メモリパフォーマンスが悪い ( )Oakforest-PACS では選択不可 2017/1/30 FrontISTR 研究会 8

9 クラスタモード : Quadrant/Hemisphere 仮想的に 4 つまたは 2 つの象限に分割 同じ象限に Tile のディレクトリとメモリが配置されるよう, アドレスをハッシュする All to All よりはメモリパフォーマンスが良い 図 5 Quadrant の概念 2017/1/30 FrontISTR 研究会 9

10 クラスタモード : SNC-4/SNC-2 4 つまたは 2 つの NUMA に分割 Intel が公開しているシェルスクリプトなどを利 して, 各 NUMA に所属してるタイルディレクトリとメモリを bind させることが可能 メモリパフォーマンスが最も良い 図 6 SNC-4 の概念 ( )Oakforest-PACS では選択不可 2017/1/30 FrontISTR 研究会 10

11 ソースコード命令による MCDRAM 利 法 Intel の資料 [3] より抜粋 [3] Shuo Li, Karthik Raman, Ruchira Sasanka, Andrey Semin, Enhancing Application Performance using Heterogeneous Memory Architectures on the Many-core Platform, 2016, 2017/1/30 FrontISTR 研究会 11

12 DDR4 と MCDRAM の特徴 WciL: Worst case interrupt Latency 図 7 要求メモリサイズとレイテンシの関係 [3] [3] Shuo Li, Karthik Raman, Ruchira Sasanka, Andrey Semin, Enhancing Application Performance using Heterogeneous Memory Architectures on the Many-core Platform, 2016, 2017/1/30 FrontISTR 研究会 12

13 KNL 搭載機材の性能と FX10, FX100 との 較 搭載 CPU 表 1 KNL 搭載機材及び FX10, FX100 との 較 Oakforest-PACS FX10 FX100 Xeon Phi Processor 7250 SPARC64 IXfx SPARC64 XIfx 理論演算性能 [Gflops] コア数 スレッド数 動作周波数 [GHz] Flops/Clock 32 [4] 8 [5] 16 [6] メモリサイズ [GB] DDR4 MCDRAM 理論メモリ速度 [GB/s] Hardware Byte/Flop 理論演算性能 = コア数 動作周波数 Flops/Clock Hardware Byte/Flop = 理論メモリ速度 / 理論演算性能 [4] David Kanter, Knights Landing Details, [5] 東京 学情報基盤センタースーパーコンピュ ティング部, 第 2 章 FX10 スーパーコンピュータシステムについて. [6] 富 通株式会社次世代テクニカルコンピューティング開発本部, FUJITSU Supercomputer PRIMEHPC FX100 次世代技術への進化, /1/30 FrontISTR 研究会 13

14 STREAM triad による性能評価 2017/1/30 FrontISTR 研究会 14

15 STREAM triad による性能評価の概要 KNL 機材の性能評価をする上で,STREAM triad 測定を実施 さが STREAM_ARRAY_SIZE の double 型 (8[Byte]) 配列 a, b, c と double 型変数 scalar で積和演算を い, メモリスループットを測定するプログラム 表 2 STREAM triad のカーネル部分 #pragma omp parallel for for (j=0; j<stream_array_size; j++) a[j] = b[j]+scalar*c[j]; #endif カーネル部分より STREAM triad の要求 Byte/Flop を求める double 型 scalar はレジスタに乗ることが期待できるので double 型配列 a, b, c と演算 (Flop) が 2 つという事実に注 すれば良い STREAM triad の要求 Byte/Flop =![#$%&] ) * = *, * = /1/30 FrontISTR 研究会 15

16 STREAM triad の対ピーク性能, 演算速度限界値 STREAM triad の要求 Byte/Flop が 12.0 より STREAM triad の対ピーク性能 [%] = Hardware Byte/Flop STREAM triad の要求 Byte/Flop STREAM triad の演算速度限界値 [Gflops] 100 = Hardware Byte/Flop = 理論演算性能 Gflops STREAM triad の対ピーク性能 [%] / 100 表 3 KNL 搭載機材と, FX10, FX100 の 較 (STREAM triad の対ピーク性能, 演算速度限界値 ) Oakforest-PACS DDR4 MCDRAM FX10 FX100 理論演算性能 [Gflops] 理論メモリ速度 [GB/s] Hardware Byte/Flop STREAM triadの対ピーク性能 [%] STREAM triadの演算速度限界値 [Gflops] /1/30 FrontISTR 研究会 16

17 Oakforest-PACS の性能評価条件 使 コンパイラ : Intel C++ compiler Intel から発表されている STREAM 計測最適化条件 [7] を遵守 測定におけるコンパイルオプション -mcmodel medium -shared-intel -O3 -xmic-avx512 -DSTREAM_ARRAY_SIZE= DOFFSET=0 -DNTIMES=10 -qopenmp -qopt-streaming-stores always lmemkind stream.c ソースコード内の malloc 関数を hbw_malloc 関数へ変更 メモリスループット評価に関して, スレッド数と要求メモリサイズを変化させる [7] Karthik Raman, Optimizing Memory Bandwidth in Knights Landing on Stream Triad, /1/30 FrontISTR 研究会 17

18 Oakforest-PACS のメモリスループット評価 1 Best Rate[GB/s] 要求メモリサイズは3[GB] で固定 CACHE-QUADRANT CACHE-SNC FLAT(MCDRAM)-QUADRANT FLAT(MCDRAM)-SNC Number of Threads 図 8 thread 数を変化させたときの STREAM triad 実 結果 FLAT(MCDRAM)-QUADRANT において Intel 側が主張する MCDRAM の実性能値 [7] ( [GB/s]) に近い値が確認できた [7] Karthik Raman, Optimizing Memory Bandwidth in Knights Landing on Stream Triad, /1/30 FrontISTR 研究会 18

19 STREAM triad 実 時の演算速度算出 次に,STREAM triad 実 時の演算速度から評価を う 表 2 STREAM triad のカーネル #pragma omp parallel for for (j=0; j<stream_array_size; j++) a[j] = b[j]+scalar*c[j]; #endif 要求メモリサイズ 3[GB] STREAM_ARRAY_SIZE = 134,217,728 より 総演算数は 2 134,217,728 [Flop] 0.27[Gflop] STREAM triad 実 時の演算速度 [Gflops] 0.27[Gflop] &JKLM&N %OP&[M&Q] 2017/1/30 FrontISTR 研究会 19

20 STREAM triad 実 時の演算速度評価 Performance[Gflops] FLAT (MCDRAM) -QUADRANT Number of Threads 図 9 STREAM triad 実 時の演算速度表 4 OFPの理論演算性能と STREAM triadの対ピーク性能, 演算速度限界値 Oakforest-PACS DDR4 MCDRAM 理論演算性能 [Gflops] STREAM triadの対ピーク性能 [%] STREAM triadの演算速度限界値 [Gflops] 最 で 40.27[Gflops] (64 thread) 演算速度限界値の約 93.7% 2017/1/30 FrontISTR 研究会 20

21 Oakforest-PACS のメモリスループット評価 Thread 数は 68 で固定 Best Rate[GB/s] CACHE-QUADRANT CACHE-SNC4 FLAT(MCDRAM)-QUADRANT FLAT(MCDRAM)-SNC MB 16GB 16GB 3K 30K 300K 3M 30M 300M 3G 30G Demand Memory Size[Byte] 図 10 要求メモリサイズを変化させたときの STREAM Triad 実 結果 要求メモリサイズが 16GB 以下の場合 要求メモリサイズに 例して測定値が上昇 16GBを超えた場合 DDR4が使われるため, 性能が下降 2017/1/30 FrontISTR 研究会 21

22 Oakforest-PACS のメモリスループット評価 3 1, Thread 数は 68 で固定 連続 定ストライド (1) Best Rate[GB/s] 間接 , ,000 24,300, ,000,000 21,870,000,000 Demand Memory Size[Byte] 34MB 16GB 図 11 配列要素へのアクセス 法を変化させたときのメモリ性能 配列要素へのアクセスが連続でない時, メモリ性能値がさがる 2017/1/30 FrontISTR 研究会 22

23 Oakforest-PACS のメモリスループット評価 4 Best Rate[GB/s] Thread 数は 68 で固定要求メモリサイズは 3GB CACHE-QUADRANT CACHE-SNC4 FLAT(MCDRAM)-QUADRANT FLAT(MCDRAM)-SNC ,024 Number of Stride 図 12 ストライド数を変化させたときのメモリ性能 ストライド数を 32 まで増加した時, メモリ性能値は直線的に下がる 2017/1/30 FrontISTR 研究会 23

24 FrontISTR による性能評価 2017/1/30 FrontISTR 研究会 24

25 FrontISTR 実 例 : Solid-100 使 コンパイラ : Intel Fortran compiler コンパイルオプション -O2 qopenmp -xmic-avx512 -lmemkind -mcmodel=medium -shared-intel -qopt-streaming-stores=always 節点数 : 1,000,000 要素数 : 970,299 由度数 : 3,000,000 共役勾配法の前処理 : 対 スケーリング必要メモリサイズ ( 全体 ): 2.15[GB] 必要メモリサイズ ( ベクトル ): 22.9[MB] (MPI 前処理 メッシュファイル分割 ) ファイル 図 13 Solid-100 のメッシュ図 全体剛性 列の作成 クラスタモードは QUADRANT に固定 FrontISTR 実 の流れ solver/precond solver/matvec solver/communication 解が収束するまで反復 解析結果ファイル出 2017/1/30 FrontISTR 研究会 25

26 SpMV の要求 Byte/Flop 表 5 FrontISTR のホットスポット (CSR 形式による SpMV) do blocknum = 0, numofblockperthread - 1 blockindex = blocknum * numofthread + threadnum do i = startpos(blockindex), endpos(blockindex) X1= X(3*i-2) X2= X(3*i-1) X3= X(3*i ) YV1= D(9*i-8)*X1 + D(9*i-7)*X2 + D(9*i-6)*X3 YV2= D(9*i-5)*X1 + D(9*i-4)*X2 + D(9*i-3)*X3 YV3= D(9*i-2)*X1 + D(9*i-1)*X2 + D(9*i )*X3 js= indexl(i-1) + 1 je= indexl(i ) do j= js, je in = iteml(j) X1= X(3*in- 2) カーネル部分 X2= X(3*in- 1) X3= X(3*in ) YV1= YV1 + AL(9*j- 8)*X1 + AL(9* j-7)*x2 + AL( 9*j-6)* X3 YV2= YV2 + AL(9*j- 5)*X1 + AL(9* j-4)*x2 + AL( 9*j-3)* X3 YV3= YV3 + AL(9*j- 2)*X1 + AL(9* j-1)*x2 + AL( 9*j )* X3 enddo js= indexu(i- 1) + 1 je= indexu(i ) do j= js, je in = itemu(j) X1= X(3*in- 2) X2= X(3*in- 1) X3= X(3*in ) YV1= YV1 + AU(9*j- 8)*X1 + AU(9* j-7)*x2 + AU( 9*j-6)* X3 YV2= YV2 + AU(9*j- 5)*X1 + AU(9* j-4)*x2 + AU( 9*j-3)* X3 YV3= YV3 + AU(9*j- 2)*X1 + AU(9* j-1)*x2 + AU( 9*j )* X3 enddo Y(3*i-2)= YV1 Y(3*i-1)= YV2 Y(3*i )= YV3 enddo enddo SpMVの要求 Byte/Flop [8] FrontISTR 研究会, HEC-MW における重要な変数のデータ格納形式, SpMV の要求 Byte/Flop を求める 図 14 係数 列に関する内部表現 [8] X, Y の要素数は 3NP indexl, indexu の要素数は NP キャッシュに乗ることが期待できる 疎 列の ゼロ要素数に合わせて きな配列となる AL, AU, iteml, itemu について考慮 float 型配列 (8byte): Integer 型配列 (4byte): AL, AU iteml, itemu 2017/1/30 FrontISTR 研究会 26 =! #$%& RS, #$%& T * * R * = UV T! 4.22

27 FrontISTR の対ピーク性能, 演算速度限界値 SpMV の要求 Byte/Flop が 4.22 より ForntISTR の対ピーク性能 [%] = Hardware Byte/Flop Hardware Byte/Flop 100 = 100 SpMVの要求 Byte/Flop 4.22 FrontISTR の演算速度限界値 [Gflops] = 理論演算性能 Gflops FrontISTR の対ピーク性能 [%] / 100 表 6 KNL 搭載機材と, FX10, FX100 の 較 (FrontISTR の対ピーク性能, 演算速度限界値 ) Oakforest-PACS DDR4 MCDRAM FX10 FX100 理論演算性能 [Gflops] 理論メモリ速度 [GB/s] Hardware Byte/Flop FrontISTRの対ピーク性能 [%] FrontISTRの演算速度限界値 [Gflops] /1/30 FrontISTR 研究会 27

28 FrontISTR DDR4 と MCDRAM の性能 較 FLAT (MCDRAM) FLAT (DDR4) FLAT (MCDRAM) FLAT (DDR4) Speed up 8 Time[s] 各メモリの逐次計算時間が基準 Number of Process Number of Process 図 15-1 MPI 並列時のスケーリング性能 図 15-2 MPI 並列時の実 時間 DDR4 使 時は,32 並列で性能向上が頭打ち 2017/1/30 FrontISTR 研究会 28

29 FrontISTR メモリモードの性能 較 FLAT (MCDRAM) CACHE FLAT (MCDRAM) CACHE Speed up 8 Time[s] FLAT (MCDRAM) での逐次計算時間が基準 Number of Process Number of Process 図 16-1 MPI 並列時のスケーリング性能 図 16-2 MPI 並列時の実 時間 わずかな差で,FLAT (MCDRAM) のほうが性能が い 2017/1/30 FrontISTR 研究会 29

30 FrontISTR 実 時の演算性能算出 FrontISTR の総演算数は疎 列の ゼロ要素数 (Num of NZ) と CG 法の反復回数 (iteration) で決まる Solid-100 では Num of NZ 256,172,328 iteration 100に固定 ゼロ要素 1 つに 2flop( 積和演算 ) が われるので 総演算数 [flop] = (Num of NZ) 2 iteration よって FrontISTR 実 時の演算性能 [Gflops] = (Num of NZ) 2 iteration MZJ[&\ PK%[&Q %OP&[M&Q] 2017/1/30 FrontISTR 研究会 30

31 FrontISTR 実 時の演算性能 Performance[Gflops] 表 7 OFP の理論演算性能と FrontISTR の対ピーク性能, 演算速度限界値 FLAT (MCDRAM) FLAT (DDR4) CACHE Number of Process Oakforst-PACS 理論演算性能 [Gflops] DDR4 MCDRAM FrontISTRの対ピーク性能 [%] FrontISTRの演算速度限界 [Gflops] 図 17 MPI 並列時の演算性能 FrontISTR の演算速度限界値に対し, 最 で FLAT (DDR4): 約 65.18% FLAT (MCDRAM): 約 40.48% 出ている 2017/1/30 FrontISTR 研究会 31

32 まとめ KNL はハードウェアとして DDR4,MCDRAM を搭載した CPU メモリモード, クラスタモードという合計 9 通りの使い がある メモリモード (DDR4,MCDRAM の使い を決める ) クラスタモード ( タイル, メモリ間の通信 法を決める ) Flat Cache Hybrid ( ) All to All ( ) Quadrant/Hemisphere SNC(Sub-NUMA Clustering) ( ) ( )Oakforest-PACS では選択不可 STREAM による Oakforest-PACS 性能測定 Flat(MCDRAM)-Quadrant という条件下で intel の主張する 475[GB/s] [7] に近い性能を確認した メモリアクセスが間接参照のとき, 連続参照に べ約 10[%] 以下まで性能が下がることを確認した FrontISTR による Oakforest-PACS 性能測定 DDR4 では 32 並列で性能向上が頭打ちなのに対し,MCDRAM は 68 並列まで性能向上できることを確認した MCDRAM 使 時, FrontISTR の演算速度限界値が 116[Gflops] なのに対し, 実性能値として 46.96[Gflops] という結果がでた 割合として約 40.48% の性能値である [7] Karthik Raman, Optimizing Memory Bandwidth in Knights Landing on Stream Triad, /1/30 FrontISTR 研究会 32

33 参考 献 [1] Avinash Sodani, Intel Xeon Phi Processor Knights Landing Architectural Overview, [2] Avinash Sodani, Knights Landing (KNL):2nd Generation Intel Xeon Phi Processor, [3] Shuo Li, Karthik Raman, Ruchira Sasanka, Andrey Semin, Enhancing Application Performance using Heterogeneous Memory Architectures on the Many-core Platform, [4] David Kanter, Knights Landing Details, 2014, ( accessed). [5] 東京 学情報基盤センタースーパーコンピュ ティング部, 第 2 章 FX10 スーパーコンピュータシステムについて, ( accessed). [6] 富 通株式会社次世代テクニカルコンピューティング開発本部, FUJITSU Supercomputer PRIMEHPC FX100 次世代技術への進化, [7] Karthik Raman, Optimizing Memory Bandwidth in Knights Landing on Stream Triad, [8] FrontISTR 研究会, HEC-MW における重要な変数のデータ格納形式, , (accessed ). [9] James Jeffers,James Reinders, Avinash Sodani, Intel Xeon Phi Processor High Performance Programming, Second Edition: Knights Landing Edition,Morgan Kaufmann, /1/30 FrontISTR 研究会 33