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

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

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

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

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

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

メニーコアプロセッサにおける最適化 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

[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

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

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

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

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

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

性能評価 • 性能評価にあたっては，並列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

• http://www.ffte.jp/

並列一次元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

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

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

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

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

背景 • 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

方針 • 並列三次元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

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

三次元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

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

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

一次元分割の場合の通信時間 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

二次元分割の場合の通信時間 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

一次元分割と二次元分割の場合の 通信時間の比較（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

一次元分割と二次元分割の場合の 通信時間の比較（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

性能評価 • 性能評価にあたっては，二次元分割を行った並列三次 元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

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

考察（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

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

並列三次元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

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

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

背景 • 近年，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

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

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

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

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

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

性能評価 • 性能評価にあたっては，以下の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

• http://www.ffte.jp/
• http://www.fftw.org/

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

並列三次元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

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

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

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

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

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