第7回 配信講義 計算科学技術特論B(2024)

521 Views

May 10, 24

スライド概要

第7回 5月30日 大規模系での高速フーリエ変換2
第6回に引き続いて、大規模系での高速フーリエ変換について講義する。 特に自動チューニング、二次元分割を用いた並列三次元FFTアルゴリズム、GPUクラスタにおける並列三次元FFTを紹介する。

シェア

またはPlayer版

埋め込む »CMSなどでJSが使えない場合

関連スライド

各ページのテキスト
1.

大規模系での高速フーリエ変換2 高橋大介 [email protected] 筑波大学計算科学研究センター 2024/5/30 計算科学技術特論B 1

2.

講義内容 • メニーコアクラスタ上の並列FFTにおける 通信隠蔽の自動チューニング • 二次元分割を用いた並列三次元FFT アルゴリズム • GPUクラスタにおける並列三次元FFT 2024/5/30 計算科学技術特論B 2

3.

メニーコアクラスタ上の並列FFTに おける通信隠蔽の自動チューニング 2024/5/30 計算科学技術特論B 3

4.

背景 • 高速フーリエ変換(fast Fourier transform,以下FFT)は科 学技術計算において今日広く用いられているアルゴリズム. • 分散メモリ型並列計算機においてチューニングを行う際に, 最適な性能パラメータはプロセッサのアーキテクチャ,ノー ド間を結合するネットワーク,そして問題サイズなどに依存 するため,これらのパラメータをその都度手動でチューニン グすることは困難になりつつある. • そこで,自動チューニングを適用したFFTライブラリとして FFTWやSPIRALなどが提案されている. • また,並列FFTにおける演算と通信のオーバーラップ手法 が提案されている[Doi and Negishi 2010]. • 並列一次元FFTにおいて通信隠蔽のパラメータを自動チュ ーニングしメニーコアクラスタ上で性能評価を行う. 2024/5/30 計算科学技術特論B 4

5.

Six-Step FFT [Bailey 90]に基づいた 並列一次元FFTアルゴリズム 全対全通信 ひねり係数 の乗算 全対全通信 全対全通信 2024/5/30 計算科学技術特論B 5

6.

メニーコアプロセッサにおける最適化 COMPLEX*16 X(N1,N2),Y(N2,N1) !$OMP PARALLEL DO COLLAPSE(2) PRIVATE(I,J,JJ) DO II=1,N1,NB DO JJ=1,N2,NB DO I=II,MIN(II+NB-1,N1) DO J=JJ,MIN(JJ+NB-1,N2) Y(J,I)=X(I,J) END DO END DO 最外側ループ長を大きくするために, END DO OpenMPのcollapse節を用いて二重 END DO ループを一重化する. !$OMP PARALLEL DO DO I=1,N1 CALL IN_CACHE_FFT(Y(1,I),N2) END DO … 2024/5/30 計算科学技術特論B 6

7.

[Idomura et al. 2014]の手法に基づく 演算と通信のオーバーラップの記述例 CALL MPI_INIT(IERR) … !$OMP PARALLEL !$OMP MASTER オーバーラップするMPI通信 !$OMP END MASTER ←同期なし !$OMP DO SCHEDULE(DYNAMIC) DO I=1,N ←マスタースレッドで通信を実行 (通信が早く終われば演算に参加可能) ←マスタースレッド以外のスレッドで 演算を実行 オーバーラップする演算 END DO ←暗黙の同期 !$OMP DO DO I=1,N 通信結果を使用する演算 END DO !$OMP END PARALLEL CALL MPI_FINALIZE(IERR) STOP END 2024/5/30 計算科学技術特論B 7

8.

演算と通信のオーバーラップ(1/2) オーバーラップ なし オーバーラップ あり(2分割) 演算 演算 通信 演算 オーバーラップ あり(4分割) 2024/5/30 通信 通信 短縮 短縮 計算科学技術特論B 8

9.

自動チューニング手法 2024/5/30 計算科学技術特論B 9

10.

通信メッセージサイズの分割数 • 演算と通信をパイプライン方式でオーバーラップさせる場 合,通信メッセージサイズの分割数(パイプラインの段数) を大きくすれば,オーバーラップの割合が高くなる. • その一方で,通信メッセージサイズの分割数を大きくすれ ば,通信1回あたりの通信メッセージサイズが小さくなるた め,通信バンド幅も小さくなる. • また,演算と通信をオーバーラップさせる際には,通信に よってメモリバンド幅が消費される. • したがって通信メッセージサイズの分割数には最適な値が 存在すると考えられる. • 通信メッセージサイズの分割数をNDIVとしたとき, NDIV=1(オーバーラップなし)~16の範囲で,通信すべき 要素数がNDIVで割り切れるすべての場合について全探 索を行う. 2024/5/30 計算科学技術特論B 10

11.

基底 2024/5/30 計算科学技術特論B 11

12.

ブロックサイズ • Six-step FFTアルゴリズムにおいては行列の転置が 必要になるが,この行列の転置はキャッシュブロッキン グを行うことで効率よく実行できることが知られている. • その際,最適なブロックサイズNBは,問題サイズおよ びキャッシュサイズ等に依存する. • 今回の実装では,ブロックサイズNBを2のべき乗に限 定して4,8,16,32,64と変化させている. 2024/5/30 計算科学技術特論B 12

13.

性能評価 • 性能評価にあたっては,並列FFTライブラリであるFFTE 6.2alpha (http://www.ffte.jp/)と,自動チューニング手法をFFTE 6.2alphaに適 用したもの,そしてFFTW 3.3.6-pl1との性能比較を行った. • 順方向FFTを1~128MPIプロセス(1ノードあたり1MPIプロセス)で連続 10回実行し,その平均の経過時間を測定した. • Xeon Phiクラスタとして,最先端共同HPC基盤施設(JCAHPC)に設置 されているOakforest-PACS(8208ノード)の128ノードを用いた. – CPU:Intel Xeon Phi 7250 (68 cores, Knights Landing 1.4 GHz) – インターコネクト:Intel Omni-Path Architecture – コンパイラ:Intel Fortran compiler 17.0.1.132 (FFTE) Intel C compiler 17.0.1.132 (FFTW) – コンパイラオプション:“-O3 -xMIC-AVX512 -qopenmp” – MPIライブラリ:Intel MPI 2017.1.132 – flat/quadrantモードおよびMCDRAMのみを用い,KMP_AFFINITY=compact – 各ノードあたりのスレッド数は64に設定 2024/5/30 計算科学技術特論B 13

14.

並列一次元FFTにおける自動チューニングの 結果(Oakforest-PACS,128ノード) FFTE 6.2alpha FFTE 6.2alpha with AT N N1 N2 NB NDIV GFlops N1 N2 NB NDIV GFlops 16M 4K 4K 32 4 57.8 4K 4K 32 1 109.4 64M 8K 8K 32 4 116.9 8K 8K 16 2 154.8 256M 16K 16K 32 4 73.3 8K 32K 16 1 513.8 1G 32K 32K 32 4 541.7 32K 32K 64 4 554.9 4G 64K 64K 32 4 217.0 64K 64K 32 16 516.5 2024/5/30 計算科学技術特論B 14

15.

700 600 500 400 300 200 100 0 32 30 28 26 24 FFTE 6.2alpha (no overlap) FFTE 6.2alpha (NDIV=4) FFTE 6.2alpha with AT FFTW 3.3.6pl1 22 20 GFlops 並列一次元FFTの性能 (Oakforest-PACS,128ノード) Length of transform (log_2 N) 2024/5/30 計算科学技術特論B 15

16.

2500 MPI_Alltoall 2000 1500 1000 500 4M 64 K 51 2K 8K 12 8 1K 0 16 Bandwidth (MB/sec) 全対全通信の性能 (Oakforest-PACS,128ノード) Message size (bytes) 2024/5/30 計算科学技術特論B 16

17.

FFTE 6.2alpha(no overlap)の実行時間の内訳 (Oakforest-PACS, N=2^26×ノード数) 3 Communication Time (sec) 2.5 Computation 2 1.5 1 0.5 0 1 2024/5/30 2 4 8 16 32 64 128 Number of nodes 計算科学技術特論B 17

18.

二次元分割を用いた並列三次元 FFTアルゴリズム 2024/5/30 計算科学技術特論B 18

19.

背景 • 2021年11月のTop500リストにおいて,2システムが 100PFlopsの大台を突破している. – Fugaku: 442,010.0 TFlops (7,630,848 Cores) – Summit: 148,600.0 TFlops (2,414,592 Cores) • 第1位のFugakuは,700万コアを超える規模になって いる. 2024/5/30 計算科学技術特論B 19

20.

方針 • 並列三次元FFTにおける典型的な配列の分散方法 – 三次元(x,y,z方向)のうちの一次元のみ(例えばz方向) のみを分割して配列を格納. – MPIプロセスが1万個の場合,z方向のデータ数が1万点 以上でなければならず,三次元FFTの問題サイズに制約. • x,y,z方向に三次元分割する方法が提案されている [Eleftheriou et al. ’05, Fang et al. ’07]. – 各方向のFFTを行う都度,全対全通信が必要. • 二次元分割を行うことで全対全通信の回数を減らしつ つ,比較的少ないデータ数でも高いスケーラビリティを 得る. 2024/5/30 計算科学技術特論B 20

21.

z方向に一次元ブロック分割した 場合の並列三次元FFT 1. x方向FFT 2. y方向FFT z z z x y 3. z方向FFT x y y x 各プロセッサでslab形状に分割 2024/5/30 計算科学技術特論B 21

22.

三次元FFTの超並列化 • 並列アプリケーションプログラムのいくつかに おいては,三次元FFTが律速になっている. • x,y,zのうちz方向のみに一次元分割した場合, 超並列化は不可能. – 1,024×1,024×1,024点FFTを2,048プロセスで 分割できない(1,024プロセスまでは分割可能) • y,zの二次元分割で対応する. – 1,024×1,024×1,024点FFTが1,048,576 (=1,024×1,024)プロセスまで分割可能になる. 2024/5/30 計算科学技術特論B 22

23.

y,z方向に二次元ブロック分割 した場合の並列三次元FFT 1. x方向FFT 2. y方向FFT z z z x y 3. z方向FFT x y y x 各プロセッサで直方体形状に分割 2024/5/30 計算科学技術特論B 23

24.

二次元分割による並列三次元FFTの実装 P×Q 2024/5/30 計算科学技術特論B 24

25.

一次元分割の場合の通信時間 P×Q N W N /(PQ) 2 PQ − 1  16 N   T1dim = ( PQ − 1) L + 2 ( PQ) ⋅ W   16 N ≈ PQ ⋅ L + (sec) PQ ⋅ W 2024/5/30 計算科学技術特論B 25

26.

二次元分割の場合の通信時間 N /( P 2Q ) N /( PQ 2 ) Q −1   16 N  16 N    + (Q − 1) L + T2 dim = ( P − 1) L + 2 2 P Q ⋅W  PQ ⋅ W    32 N ≈ ( P + Q) ⋅ L + (sec) PQ ⋅ W 2024/5/30 計算科学技術特論B 26

27.

一次元分割と二次元分割の場合の 通信時間の比較(1/2) 16 N T1dim ≈ PQ ⋅ L + PQ ⋅ W 32 N T2 dim ≈ ( P + Q) ⋅ L + PQ ⋅ W P×Q 2024/5/30 計算科学技術特論B 27

28.

一次元分割と二次元分割の場合の 通信時間の比較(2/2) • 二次元分割の通信時間が一次元分割の通信時間より も少なくなる条件を求める. 32 N 16 N < PQ ⋅ L + ( P + Q) ⋅ L + PQ ⋅ W PQ ⋅ W を解くと, ( LW ⋅ PQ)( PQ − P − Q) N< 16 9 −5 = 10 W 10 L = • 例えば, (sec), (Byte/s),P = Q = 64 10 < 10 N を上の式に代入すると, の範囲では二次元 分割の通信時間が一次元分割に比べて少なくなる. 2024/5/30 計算科学技術特論B 28

29.

性能評価 • 性能評価にあたっては,二次元分割を行った並列三次 元FFTと,一次元分割を行った並列三次元FFTの性能 比較を行った. 3 3 3 3 N = 32 , 64 , 128 , 256 • Strong Scalingとして 点の順方向 FFTを1~4,096MPIプロセスで連続10回実行し,その 平均の経過時間を測定した. • 評価環境 – T2K筑波システムの256ノード(4,096コア)を使用 – flat MPI(1core当たり1MPIプロセス) – MPIライブラリ:MVAPICH 1.2.0 – Intel Fortran Compiler 10.1 – コンパイルオプション:”ifort –O3 –xO”(SSE3ベクトル命令) 2024/5/30 計算科学技術特論B 29

30.

1000 二次元分割を行ったvolumetric 並列三次元FFTの性能 GFLOPS N=32^3 100 N=64^3 10 N=128^3 1 N=256^3 96 40 24 10 6 25 64 16 4 1 0.1 Number of cores 2024/5/30 計算科学技術特論B 30

31.

考察(1/2) • N = 32 3 点FFTでは良好なスケーラビリティが得られて いない. • これは問題サイズが小さい(データサイズ:1MB)こと から,全対全通信が全実行時間のほとんどを占めて いるからであると考えられる. 3 N = 256 • それに対して, 点FFT(データサイズ:512MB) では4,096コアまで性能が向上していることが分かる. – 4,096コアにおける性能は約401.3 GFlops (理論ピーク性能の約1.1%) – 全対全通信を除いたカーネル部分の性能は約10.07 TFlops (理論ピーク性能の約26.7%) 2024/5/30 計算科学技術特論B 31

32.

GFLOPS 1000 256^3点FFTにおける一次元分割と 二次元分割の性能比較 一次元 分割 二次元 分割 100 10 1 96 40 24 10 6 25 64 16 4 1 0.1 Number of cores 2024/5/30 計算科学技術特論B 32

33.

並列三次元FFTの実行時間の内訳 (256cores, 256^3点FFT) 0.06 Time (sec) 0.05 演算時間 通信時間 0.04 0.03 0.02 0.01 0 一次元分割 2024/5/30 二次元分割 計算科学技術特論B 33

34.

考察(2/2) • 64コア以下の場合には,通信量の少ない一次元分割 が二次元分割よりも性能が高くなっている. • 128コア以上では通信時間を少なくできる二次元分割 が一次元分割よりも性能が高くなっていることが分かる. • 二次元分割を行った場合でも,4,096コアにおいては 96%以上が通信時間に費やされている. – 全対全通信において各プロセッサが一度に送る通信量が わずか1KBとなるため,通信時間においてレイテンシが 支配的になるためであると考えられる. • 全対全通信にMPI_Alltoall関数を使わずに,より低レベ ルな通信関数を用いて,レイテンシを削減する工夫が 必要. 2024/5/30 計算科学技術特論B 34

35.

GPUクラスタにおける 並列三次元FFT 2024/5/30 計算科学技術特論B 35

36.

背景 • 近年,GPU(Graphics Processing Unit)の高い演算 性能とメモリバンド幅に着目し,これを様々なHPCアプ リケーションに適用する試みが行われている. • また,GPUを搭載した計算ノードを多数接続したGPU クラスタも普及が進んでおり,2021年11月のTOP500 リストではNVIDIA V100 GPUを搭載したSummitが第 2位にランクされている. • これまでにGPUクラスタにおける並列三次元FFTの実 現は行われている[Chen et al. 2010, Nukada et al. 2012]が,一次元分割のみサポートされており,二次 元分割はサポートされていない. 2024/5/30 計算科学技術特論B 36

37.

方針 • CPU版とGPU版を同一インターフェースとするため, 入力データおよび出力データはホストメモリに格納す る. – FFTライブラリが呼び出された際に,ホストメモリからデバイ スメモリに転送し,FFTライブラリの終了時にデバイスメモリ からホストメモリに転送する. • 計算可能な問題サイズはGPUのデバイスメモリの容 量が限度になる. – ホストメモリのデータを分割してデバイスメモリに転送しなが らFFT計算を行うことも可能であるが,今回の実装ではそこ まで行わないこととする. 2024/5/30 計算科学技術特論B 37

38.

並列三次元FFTアルゴリズム 全対全通信 転置 全対全通信 転置 38

39.

GPUクラスタにおける並列三次元FFT(1/2) • GPUクラスタにおいて並列三次元FFTを行う際には, 全対全通信が2回行われる. • 計算時間の大部分が全対全通信によって占められる ことになる. • CPUとGPU間を接続するインターフェースであるPCI Expressバスの理論ピークバンド幅はPCI Express Gen 2 x 16レーンの場合には一方向あたり8GB/sec. • CPUとGPU間のデータ転送量をできるだけ削減するこ とが重要になる. – CPUとGPU間のデータ転送はFFTの開始前と終了後にそ れぞれ1回のみ行う. – 行列の転置はGPU内で行う. 2024/5/30 計算科学技術特論B 39

40.

GPUクラスタにおける並列三次元FFT(2/2) • GPU上のメモリをMPIにより転送する場合,以下の手 順で行う必要がある. 1. GPU上のデバイスメモリからCPU上のホストメモリへデー タをコピーする. 2. MPIの通信関数を用いて転送する. 3. CPU上のホストメモリからGPU上のデバイスメモリにコピー する. • この場合,CPUとGPUのデータ転送を行っている間は MPIの通信が行われないという問題がある. • そこで,CPUとGPU間のデータ転送とノード間のMPI 通信をパイプライン化してオーバーラップさせることが できるMPIライブラリであるMVAPICH2を用いた. 2024/5/30 計算科学技術特論B 40

41.

MPI + CUDAでの通信 • 通常のMPIを用いたGPU間の通信 At Sender: cudaMemcpy(sbuf, s_device, …); ・cudaMemcpyを行っている間 はMPIの通信が行われない. MPI_Send(sbuf, size, …); ・メモリをブロックで分割し, At Receiver: CUDAとMPIの転送をオーバ MPI_Recv(rbuf, size, …); ーラップさせることも可能. cudaMemcpy(r_device, rbuf, …); →プログラムが複雑になる. • MVAPICH2-GPUを用いたGPU間の通信 At Sender: ・デバイスメモリのアドレスを MPI_Send(s_device, size, …); At Receiver: MPI_Recv(r_device, size, …); 2024/5/30 直接MPI関数に渡すことが可能. ・CUDAとMPIの転送のオーバー ラップをMPIライブラリ内で行う. 計算科学技術特論B 41

42.

性能評価 • 性能評価にあたっては,以下のFFTライブラリについて性能比較を行った. – FFTE 6.0(http://www.ffte.jp/,GPUを使用) – FFTE 6.0(http://www.ffte.jp/,CPUを使用) – FFTW 3.3.3(http://www.fftw.org/,CPUを使用) • 順方向FFTを1~256MPIプロセス(1ノードあたり4MPIプロセス)で連続 10回実行し,その平均の経過時間を測定した. • HA-PACSベースクラスタ(268ノード,4288コア,1072GPU)の うち,1~64ノードを使用した. – 各ノードにIntel Xeon E5-2670(Sandy Bridge-EP 2.6GHz)が2ソケット, NVIDIA Tesla M2090が4基 – ノード間はInfiniBand QDR(2レール)で接続 – MPIライブラリ:MVAPICH2 2.0b – PGI CUDA Fortran Compiler 14.2 + CUDA 5.5 + CUFFT – コンパイラオプション:“pgf90 -fast -Mcuda=cc2x,cuda5.5”(FFTE 6.0,GPU), “pgf90 –fast -mp”(FFTE 6.0,CPU),”pgcc -fast”(FFTW 3.3.3) 2024/5/30 計算科学技術特論B 42

43.

HA-PACSベースクラスタのノード構成 1GPUあたり 1MPIプロセス を立ち上げる 2024/5/30 計算科学技術特論B 43

44.

並列三次元FFTの性能 (HA-PACS,N=256×256×512×MPIプロセス数) GFlops 1000 FFTE 6.0 (GPU) 100 FFTE 6.0 (CPU) 10 FFTW 3.3.3 (CPU) 256 128 64 32 16 8 4 2 1 1 Number of MPI processes 2024/5/30 計算科学技術特論B 44

45.

FFTE 6.0(GPU版)の並列三次元FFTの実行時間の内訳 (HA-PACS,N=256×256×512×MPIプロセス数) 3 通信時間 Time (sec) 2.5 2 PCIe転送時 間 1.5 演算時間 1 0.5 8 6 12 25 64 32 16 8 4 2 1 0 Number of MPI processes 2024/5/30 計算科学技術特論B 45

46.

FFTE 6.0(CPU版)の並列三次元FFTの実行時間の内訳 (HA-PACS,N=256×256×512×MPIプロセス数) 3 通信時間 Time (sec) 2.5 演算時間 2 1.5 1 0.5 8 6 12 25 64 32 16 8 4 2 1 0 Number of MPI processes 2024/5/30 計算科学技術特論B 46

47.

2M 12 K 8K 25 6K 51 2K 1M 64 K 32 16 K GPU-GPU (with MVAPICH2GPU) GPU-GPU (without MVAPICH2GPU) CPU-CPU 8K 4K Bandwidth (MB/sec) 800 700 600 500 400 300 200 100 0 全対全通信の性能 (HA-PACS 64ノード, 256MPIプロセス) Message Size (bytes) 2024/5/30 計算科学技術特論B 47

48.

まとめ(1/2) • 物質科学の実アプリケーションにおいて使われることが 多い,高速フーリエ変換(FFT)について紹介した. • これまで並列FFTで行われてきた自動チューニングで は,基数の選択や組み合わせ,そしてメモリアクセスの 最適化など,主にノード内の演算性能だけが考慮され てきた. • ノード内の演算性能だけではなく,通信の隠蔽において も自動チューニングが必要になる. • 今後,並列スーパーコンピュータの規模が大きくなるに 従って、FFTの効率を向上させることは簡単ではない. – 二次元分割や三次元分割が必要がある. 2024/5/30 計算科学技術特論B 48

49.

まとめ(2/2) • GPUを用いた場合にはCPUに比べて演算時間が短縮 される一方で,全実行時間における通信時間の割合が 増大する. – HA-PACSベースクラスタの64ノード,256MPIプロセスを用い た場合,2048^3点FFTにおいて実行時間の約70%が全対全 通信で占められている. • MPIライブラリであるMVAPICH2の新機能 (MVAPICH2-GPU)を用いることで,PCIe転送とノード 間通信をオーバーラップさせた際のプログラミングが容 易になるとともに通信性能も向上した. 2024/5/30 計算科学技術特論B 49