Microsoft PowerPoint - CAEworkshop_ _01.ver1.3

Size: px

Start display at page:

Download "Microsoft PowerPoint - CAEworkshop_ _01.ver1.3"

しほこやたけ
6 years ago
Views:

1 GPU メニーコアにおける OpenFOAM の高度化支援紹介第 1 回 CAE ワークショップ流体構造解析アプリケーションを中心に 2017 年 12 月 6 日秋葉原 UDX Gallery NEXT 山岸孝輝井上義昭青柳哲雄浅見曉 ( 高度情報科学技術研究機構 ) ver 1.3 1

2 outline RISTの高度化支援について GPU メニーコアについて OpenFOAMとGPU GPU 対応版 OpenFOAMの最適化新しいGPU 版 OpenFOAM 開発について OpenFOAM とメニーコア Intel KNL での基礎性能調査まとめ 2

3 RIST の高度化支援文科省委託事業 HPCI の運営代表機関として京を中核とする HPCI の高度化を支援利用者の要望にワンストップで対応様々な場面に様々なメニューを提供します OpenFOAM の高度化支援 GPU メニーコア対応を紹介 RIST の高度化支援の具体例としても 3

4 スパコンと GPU メニーコア低消費電力低コストスパコンランキング (TOP500) 上位を占めておりその割合は増加傾向暁光 (4 位 ), Oakforest-PACS(9 位 ), TSUBAME 3.0(13 位 ) など多数ランク入り利用者の視点からは現実的には使いにくいのが現状 ( 移植高速化コードの保守管理 etc.) RIST の高度化支援にて最も大きな課題の一つ 4

5 OpenFOAM と GPU メニーコア GPU:NVIDIA GPU 実装に手間 (CUDA, OpenACC, Thrust, etc.) 性能を出すのは容易ではない GPU 実装版 (RapidCFD) が公開されているメニーコア :Intel Xeon Phi KNL 移植は比較的楽 (Intel コンパイラ, MPI) 性能については試行錯誤の段階 5

6 RapidCFD 概要 GPU 対応版 OpenFOAM simflow 社が GPU に移植開発中 : Let s make RapidCFD rock solid! 特徴リリースは 3 年前最後の bug fix は 2 年前大半のソルバに対応 Github でソースコード公開 RapidCFD-dev 全ての計算を GPU に移植 (GPU-CPU 間メモリ転送コストの削減 ) 6

7 RapidCFD 性能評価事例 (1/2) 1GPU vs 1CPU simflow 社ウェブサイトから引用 [ 秒 ] 同世代の CPU 比較で 2 倍強高速 7

8 加速率 RapidCFD 性能評価事例 (2/2) 複数 GPU でのスケーリング性能 simflow 社ウェブサイトから引用 GPU 数 4GPU から性能が鈍化 8

9 設定したターゲット RapidCFD の高速化フル移植済みの環境を活用する単体性能 (1GPU) の向上 CPU と理論で 10 倍違う性能を引き出す複数 GPU でのスケーリング性能の向上大きな問題サイズへの対応実行時間の削減 TSUBAME 3.0 での試みを RIST の高度化支援の流れと共に紹介 9

10 TSUBAME 3.0 測定環境項目内容 HW TSUBAME 3.0 SGI ICE XA CPU Xeon E5-2680v4[426GF, 2.4GHz, 14core], 2socket GPU Tesla P100 with NVLink[5.3TF, 16GB], 4board ノード間接続 Intel Omni-Path 100Gb/s x4 SW C++ Icc MPI OpenMPI CUDA CUDA AP OpenFOAM RapidCFD ThrustライブラリによるGPU 実装解析ソルバ pisofoam/pimplefoam 解析データ (1) 弱スケーリング 128x128x128 cells/processor~ (2) 強スケーリング 240x130x96 cells/1-64processros 10

11 TSUBAME3.0 ノード構成 2CPU + 4GPU 構成 CPU-GPU 間は PCIe GPU 間は高速な NVLink GPUDirect RDMA ノード間転送 TSUBAME3.0 Web から引用ノード内のリソースを使いこなしつつノード間で効率的に並列化させる 11

12 解析モデル ( 弱スケーリング ) 風洞測定部 16m 2m 3m 1 プロセッサあたり 200 or 300 万格子を設定 GPU or CPU 格子数 1 128x128x 万 2 256x128x 万 4 512x128x 万 8 256x256x 万 x256x 万 200 万格子での設定 200 万格子 /GPU or CPU で一定 ( 弱スケーリング ) 理想的には性能 ( 経過時間 ) はスケール ( 一定 ) 12

13 性能計測方法メインループの前後にタイマ挿入 timer_start("runt",1,1); while (runtime.loop()) { } < メイン処理 > timer_stop("runt",1,1); Info<< "End n" << endl; return 0; 経過時間計測範囲最初はざっくりとした方法で計測目的はコスト分布の把握余計な情報を入れない CPU GPU に縛られない方法で計測比較する 13

14 CPU vs GPU(ASIS) [ 秒 ] プロセッサ単体での性能 44.4 CPU 3.3 倍 300 万格子 /1GPU or 1CPU 13.3 GPU CPU は 12 コアで flat MPI CPU:E5-2680v4 GPU:P100 同世代の CPU に比べて 3.3 倍高速 14

15 GPU 弱スケーリング性能 (ASIS) 2 種類の GPU 配置方法で計測 (1 4GPU) ノード内並列ノード間並列 X4 [ 秒 ] [ 秒 ] 万格子 /1GPU GPU 数 GPU 数 GPU のノード内並列で性能劣化 4GPU で 2 倍低速 15

16 処理プロセスごとに分類して計測 [ 秒 ] ノード内並列メインループでの分類 GPU 数 others UEqn.2 UEqn.3 UEqn.4 PISO.2 peqn.2 peqn peqn.3 peqn.4 16

17 ソースコード解析性能阻害箇所 template<> scalar sumprod(const gpulist<scalar>& f1, const gpulist<scalar>& f2) { if (f1.size() && (f1.size() == f2.size())) { thrust::device_vector<scalar> t(f1.size()); thrust::transform ( f1.begin(), f1.end(), f2.begin(), t.begin(), multiplyoperatorfunctor<scalar,scalar,scalar>() ); thrust ライブラリによる GPU 実装総和計算処理にて性能が悪化と判明 } } return thrust::reduce(t.begin(),t.end()); else { return 0.0; } 総和計算本来ならライブラリが得意とする処理 17

18 Visual Profiler による詳細分析マイクロ秒単位 1GPU 時切り出して計測カーネル API レベル cudamalloc reduce cudafree タイムラインを追う 4GPU 時 cudamalloc GPU 上のメモリ確保解放のコストが GPU 数につれて増大 reduce cudafree 18

19 総和計算 CUDA による GPU 実装目的 GPU 上の無駄なメモリ確保解放の除去総和計算の高速化実装についてインタリーブ方式最初の足し込みステップでスレッドの削減ウォープダイバージェンスの回避 Thrust::reduce(1 行 ) CUDA カーネル (260 行 ) 低データ並列性からの CPU への切り替え 19

20 CPU vs GPU(TUNED) [ 秒 ] プロセッサ単体での性能倍 4.4 倍 CPU GPU ASIS GPU TUNED 3.3 倍 4.4 倍高速に 300 万格子 /1GPU or 1CPU 1.3 倍 CPU は 12 コアで flat MPI 20

21 [ 秒 ] GPU 弱スケーリング性能 (TUNED) ノード 4GPU CPU GPU ASIS GPUTUNED 200 万格子 /1GPU GPU 数多数 GPU 設定 (4, 8, 16) にて 2 倍改善 21

22 解析モデル 2( 強スケーリング ) OpenFOAM-BenchmarkTest 240x130x 万格子 /1-64 CPU or GPU で分割 1, 2, 4, 8, 16, 32, 64 今野雅氏第 4 回 OpenFOAM ワークショップ発表資料から引用 am_161216_imano 理想的には GPU の増加に比例して高速化 22

23 GPU 強スケーリング性能 [ 秒 ] Case1 16PE から飽和 64PE からは逆スケーリング処理プロセスごとに分類して分析 Case2 GPU 数 Case の違いは乱流モデルの扱い GPU 数増加につれてホスト (CPU) デバイス (GPU) メモリ間転送コスト割合が増加 23

24 GPUDirect RDMA 転送異なるノード間で GPU メモリが直接転送 / やりとりできる東大 Reedbush, 東工大 TSUBAME 3.0 などで採用 CPU メモリ GPU メモリ GPU メモリ GPU メモリ GPU メモリ CPU メモリ CPU GPU 2 GPU 1 GPU 1 GPU 2 CPU NIC network NIC ノード 0 ノード 1 24

25 RapidCFD 改善の余地について OpenFOAM の version が古い最新の 5.0 よりも約 3 年前のもの Thrust ライブラリでの GPU 実装柔軟な最適化が難しいコンピュートケイパビリティが 3.0 に設定プロファイラ結果の解析が困難 Time(%) Time Calls Avg Min Max Name 0.47% ms us us us Foam::reduce_C_5(double*,double*,int) Time(%) Time Calls Avg Min Max Name 0.47% ms us us ms void thrust::system::cuda::detail::bulk_::detail::launch_by_value<unsigned int=0, thrust::system::cuda::detail::bulk_::detail::cuda_task<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail:: bulk_::agent<unsigned long=1>, unsigned long=0>, unsigned long=0>, thrust::system::cuda::detail::bulk_::detail::closure<thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, thrust::tuple<thrust::system::cuda::detail::bulk_::detail::cursor<unsigned int=0>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<foam::vector<double> const >>, thrust::detail::normal_iterator<thrust::device_ptr<foam::vector<double> const >>, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::detail::wrapped_function<thrust::detail::binary_transform_functor<foam::dotoperatorfunctor<foam::vector<double>, Foam::Vector<double>, double>>, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>>(unsigned long=1) CUDA Thrust 25

26 新しい GPU 版 OpenFOAM 開発へ最新の OpenFOAM をベースに開発開発の方針特徴 GPU 実装はすべて CUDA メインループ内の処理は全て GPU で処理 GPUDirect RDMA 転送に対応現在の status 代表的なソルバを CUDA で実装本発表内容を再現することから目標高速な大規模実行の実現 GPU を用いた大規模アプリに適用しうる知見の蓄積 26

27 GPU での話題まとめ GPU 対応 OpenFOAM, RapidCFD の性能評価と最適化を行った総和計算の CUDA による実装で無駄なメモリ確保解放コストを削減 GPU 単体では倍高速に (CPU 比較 ) 複数 GPU でのスケーリング性能は最大で 2 倍改善強スケーリングでの通信ボトルネック解消には GPUDirect RDMA 転送が有効新しい GPU 版 OpenFOAM 開発へ 27

28 メニーコア :Intel Xeon Phi KNL 支援前の基礎調査として RIST 所有の KNL 搭載 PC クラスタで性能評価項目内容 HW KNL 搭載 PCクラスタ CX6000M1/CX1640M1x4 ノード数 4 モード設定 flat/cache/mcdram + Quadrant/AlltoAll/SNC2/SNC4 CPU Intel Xeon Phi 7250[3TFLOPS, 1.4GHz, 68core] Intel Omni-Path SW C++/Fortran Icc MPI Intel MPI 2017 update 1 コンパイルオプション -O3 -xmic-avx512 AP OpenFOAM 2.3.1(RIST 最適化版 ) 解析ソルバ pimplefoam 解析データ強スケーリング 240x130x96 cells/1-64processros 28

29 メモリモード変更 OpenFOAM channelretau110 参考 : 京 DRAM cache MCDRAM PCG, 前処理 DIC 1 ノード内並列数の増加 ( PE) DRAM 設定時並列数を増加させると演算処理部の時間が急増スケーリング性能劣化メモリモード変更 (DRAM cache/mcdram) で 3 倍弱高速に 29

Quadrant (quad): Four virtual address spaces (one NUMA domain) Sub-NUMA-2 (snc2): Two distinct

30 クラスターモード +Hyper-threading 各種モード (quad,a2a,snc2or4/flat,cache 等 ) Hyper-threading 全て 1node HT 無 HT 2 HT 4 KNL での最速クラスターモード及び Hyper-threading 効果を検証 All-to-all (a2a): Uniform mesh interconnect Quadrant (quad): Four virtual address spaces (one NUMA domain) Sub-NUMA-2 (snc2): Two distinct NUMA domains Sub-NUMA-4 (snc4): Four distinct NUMA domains 最速 :snc2/flat+128mpi[64corex2ht] 最大で 17% のコスト削減 30

31 マルチノード +HyperThreading KNL/4node 構成 (quad, cache モード ) Mc: McDram HT:HyperThread xn( スレッド数 ) 参考 : 京 1ノード 2ノード 4ノード (HT:x2) (x2) (x4) (x2) 1-4 ノードまで HW モード HT に依存しないでスケール 31

32 KNL での話題まとめ KNL4 ノードでの OpenFOAM の性能評価を実施メモリモード変更 (DRAM cache/mcdram)3 倍弱高速にクラスターモード HT の設定は最大で 17% 影響 Quad, SNC2, HT ON(2) が良好 4 ノードまではクラスターモード HT の設定に依存しないでスケーリングする 32

33 まとめ RIST は京を中核とする HPCI の高度化を支援しています支援を通じた OpenFOAM の GPU メニーコア対応について紹介しました GPU: 高速化 (2 倍 ) と GPU 版 OpenFOAM 構築に向けた取り組みについてメニーコア : メモリ設定 HW モード HT スケーリング性能についての基礎調査結果 33

RIST の高度化支援ざっくりと評価プロセスレベル ( 秒 ) 詳細な評価と分析カーネル API レベル ( マイクロ秒 ) コードの改善へ CUDA 実装も対応 260 行 1 行 thrust::reduce(); //====dummy code===== template<class Type> Type minmagsqr(const gpulist<type>& f) { if

begin(), magsqrunaryfunctionfunctor<type,scalar>() ); typename thrust::device_vector<scalar>::iterator iter = thrust::min_element(ms.begin(), ms.end()); unsigned int position = iter - ms.

34 RIST の高度化支援ざっくりと評価プロセスレベル ( 秒 ) 詳細な評価と分析カーネル API レベル ( マイクロ秒 ) コードの改善へ CUDA 実装も対応 260 行 1 行 thrust::reduce(); //====dummy code===== template<class Type> Type minmagsqr(const gpulist<type>& f) { if (f.size()) { gpulist<scalar>& ms (f.size()); thrust::transform ( f.begin(), f.end(), ms.begin(), magsqrunaryfunctionfunctor<type,scalar>() ); typename thrust::device_vector<scalar>::iterator iter = thrust::min_element(ms.begin(), ms.end()); unsigned int position = iter - ms.begin(); return f.get(position); } else { return ptraits<type>::rootmax; } } template<class Type> scalar sumprod(const gpulist<type>& f1, const gpulist<type>& f2) { if (f1.size() && (f1.size() == f2.size())) { fapp_start("sumprod",1,2); gpulist<scalar> tmp(f1.size()); thrust::transform ( //=====dummy code===== 全て無料利用の始まりから終わりまでワンストップ窓口が対応幅広い内容に対応 RIST 高度化支援検索 34

35 謝辞本発表の内容は,HPCI システム利用研究課題流体粒子の大規模連成解析を用いた竜巻中飛散物による建物被害の検討 ( 課題番号 :hp170055), 課題代表者 : 菊池浩利 ( 清水建設 ( 株 ) 技術研究所 ) の高度化支援に基づくものです. 35

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード] 200/0/9 数値流体解析の並列効率とその GPU による高速化の試み清水建設 ( 株 ) 技術研究所 PHAM VAN PHUC ( ファムバンフック ) 流体計算時間短縮と GPU の活用の試み現 CPUとの比較によりGPU 活用の可能性現 CPU の最大利用ノード内の最大計算資源の利用すべてCPUコアの利用適切なアルゴリズムの利用 CPU コア性能の何倍? GPU の利用の試み