GPU メニーコアにおける OpenFOAM の高度化支援紹介 第 1 回 CAE ワークショップ 流体 構造解析アプリケーションを中心に 2017 年 12 月 6 日秋葉原 UDX Gallery NEXT 山岸孝輝井上義昭青柳哲雄浅見曉 ( 高度情報科学技術研究機構 ) ver 1.3 1
outline RISTの高度化支援について GPU メニーコアについて OpenFOAMとGPU GPU 対応版 OpenFOAMの最適化新しいGPU 版 OpenFOAM 開発について OpenFOAM とメニーコア Intel KNL での基礎性能調査 まとめ 2
RIST の高度化支援 文科省委託事業 HPCI の運営 代表機関として 京 を中核とする HPCI の高度化を支援 利用者の要望にワンストップで対応 様々な場面に 様々なメニューを提供します OpenFOAM の高度化支援 GPU メニーコア対応を紹介 RIST の高度化支援の具体例としても 3
スパコンと GPU メニーコア 低消費電力 低コスト スパコンランキング (TOP500) 上位を占めており その割合は増加傾向 暁光 (4 位 ), Oakforest-PACS(9 位 ), TSUBAME 3.0(13 位 ) など多数ランク入り 利用者の視点からは 現実的には使いにくいのが現状 ( 移植 高速化 コードの保守管理 etc.) RIST の高度化支援にて最も大きな課題の一つ 4
OpenFOAM と GPU メニーコア GPU:NVIDIA GPU 実装に手間 (CUDA, OpenACC, Thrust, etc.) 性能を出すのは容易ではない GPU 実装版 (RapidCFD) が公開されている メニーコア :Intel Xeon Phi KNL 移植は比較的楽 (Intel コンパイラ, MPI) 性能については試行錯誤の段階 5
RapidCFD 概要 GPU 対応版 OpenFOAM simflow 社が GPU に移植 開発中 : Let s make RapidCFD rock solid! 特徴 リリースは 3 年前 最後の bug fix は 2 年前 大半のソルバに対応 Github でソースコード公開 https://github.com/atizar/ RapidCFD-dev 全ての計算を GPU に移植 (GPU-CPU 間メモリ転送コストの削減 ) 6
RapidCFD 性能評価事例 (1/2) 1GPU vs 1CPU simflow 社ウェブサイトから引用 [ 秒 ] 同世代の CPU 比較で 2 倍強高速 7
加速率 RapidCFD 性能評価事例 (2/2) 複数 GPU でのスケーリング性能 simflow 社ウェブサイトから引用 GPU 数 4GPU から性能が鈍化 8
設定したターゲット RapidCFD の高速化 フル移植済みの環境を活用する 単体性能 (1GPU) の向上 CPU と理論で 10 倍違う性能を引き出す 複数 GPU でのスケーリング性能の向上 大きな問題サイズへの対応実行時間の削減 TSUBAME 3.0 での試みを RIST の高度化支援の流れと共に紹介 9
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 17.0.4 MPI OpenMPI 2.1.1 CUDA CUDA 8.0.61 AP OpenFOAM 2.3.0 RapidCFD 2.3.1+ThrustライブラリによるGPU 実装 解析ソルバ pisofoam/pimplefoam 解析データ (1) 弱スケーリング 128x128x128 cells/processor~ (2) 強スケーリング 240x130x96 cells/1-64processros 10
TSUBAME3.0 ノード構成 2CPU + 4GPU 構成 CPU-GPU 間は PCIe GPU 間は高速な NVLink GPUDirect RDMA ノード間転送 TSUBAME3.0 Web から引用 ノード内のリソースを使いこなしつつ ノード間で効率的に並列化させる 11
解析モデル ( 弱スケーリング ) 風洞測定部 16m 2m 3m 1 プロセッサあたり 200 or 300 万格子を設定 GPU or CPU 格子数 1 128x128x128 200 万 2 256x128x128 400 万 4 512x128x128 800 万 8 256x256x256 1600 万 16 512x256x256 3200 万 200 万格子での設定 200 万格子 /GPU or CPU で一定 ( 弱スケーリング ) 理想的には性能 ( 経過時間 ) はスケール ( 一定 ) 12
性能計測方法 メインループの前後にタイマ挿入 timer_start("runt",1,1); while (runtime.loop()) { } < メイン処理 > timer_stop("runt",1,1); Info<< "End n" << endl; return 0; 経過時間計測範囲 最初はざっくりとした方法で計測 目的はコスト分布の把握 余計な情報を入れない CPU GPU に縛られない方法で計測 比較する 13
CPU vs GPU(ASIS) [ 秒 ] 50 40 30 20 10 0 プロセッサ単体での性能 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
GPU 弱スケーリング性能 (ASIS) 2 種類の GPU 配置方法で計測 (1 4GPU) ノード内並列 ノード間並列 X4 [ 秒 ] [ 秒 ] 25 20 14.9 15 10.6 10 5 23.2 25 20 15 10 5 10.6 11.5 12.1 200 万格子 /1GPU 0 1 2 4 GPU 数 0 1 2 4 GPU 数 GPU のノード内並列で性能劣化 4GPU で 2 倍低速 15
処理プロセスごとに分類して計測 [ 秒 ] 25 20 15 10 5 0 ノード内並列 1.1 0.2 1.5 0.7 0.8 6.0 メインループでの分類 0.3 1.6 2.1 1.0 1.0 8.4 0.4 2.1 3.6 1.5 1.5 13.5 0.1 0.1 0.1 1 2 4 GPU 数 others UEqn.2 UEqn.3 UEqn.4 PISO.2 peqn.2 peqn peqn.3 peqn.4 16
ソースコード解析性能阻害箇所 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
Visual Profiler による詳細分析 マイクロ秒単位 1GPU 時 切り出して計測 カーネル API レベル cudamalloc reduce cudafree タイムラインを追う 4GPU 時 cudamalloc GPU 上のメモリ確保 解放のコストが GPU 数につれて増大 reduce cudafree 18
総和計算 CUDA による GPU 実装 目的 GPU 上の無駄なメモリ確保 解放の除去 総和計算の高速化 実装について インタリーブ方式 最初の足し込みステップでスレッドの削減 ウォープダイバージェンスの回避 Thrust::reduce(1 行 ) CUDA カーネル (260 行 ) 低データ並列性からの CPU への切り替え 19
CPU vs GPU(TUNED) [ 秒 ] 50.0 45.0 40.0 35.0 30.0 25.0 20.0 15.0 10.0 5.0 0.0 プロセッサ単体での性能 44.4 3.3 倍 4.4 倍 13.3 10.1 CPU GPU ASIS GPU TUNED 3.3 倍 4.4 倍高速に 300 万格子 /1GPU or 1CPU 1.3 倍 CPU は 12 コアで flat MPI 20
[ 秒 ] GPU 弱スケーリング性能 (TUNED) 40.0 35.0 30.0 25.0 20.0 15.0 10.0 5.0 34.0 30.7 31.2 31.6 32.2 24.4 24.9 23.2 14.9 13.1 13.2 12.0 10.6 9.0 7.2 1 ノード 4GPU CPU GPU ASIS GPUTUNED 200 万格子 /1GPU 0.0 1 2 4 8 16 GPU 数 多数 GPU 設定 (4, 8, 16) にて 2 倍改善 21
解析モデル 2( 強スケーリング ) https://github.com/opencae/ OpenFOAM-BenchmarkTest 240x130x96 300 万格子 /1-64 CPU or GPU で分割 1, 2, 4, 8, 16, 32, 64 今野雅氏第 4 回 OpenFOAM ワークショップ発表資料から引用 http://www.hpcioffice.jp/materials/ws_openfo am_161216_imano 理想的には GPU の増加に比例して高速化 22
GPU 強スケーリング性能 [ 秒 ] 70 60 50 40 30 20 10 0 65.1 51.4 42.3 35.1 28.7 24.1 20.219.2 21.2 16.515.4 15.9 17.7 13.4 1 2 4 8 16 32 64 Case1 16PE から飽和 64PE からは逆スケーリング 処理プロセスごとに分類して分析 Case2 GPU 数 Case の違いは乱流モデルの扱い GPU 数増加につれてホスト (CPU) デバイス (GPU) メモリ間転送コスト割合が増加 23
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
RapidCFD 改善の余地について OpenFOAM の version が古い 最新の 5.0 よりも約 3 年前のもの Thrust ライブラリでの GPU 実装 柔軟な最適化が難しい コンピュートケイパビリティが 3.0 に設定 プロファイラ結果の解析が困難 Time(%) Time Calls Avg Min Max Name 0.47% 50.343ms 21318 2.3610us 1.6960us 49.216us Foam::reduce_C_5(double*,double*,int) Time(%) Time Calls Avg Min Max Name 0.47% 50.761ms 523 97.057us 5.3120us 1.0682ms 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
新しい GPU 版 OpenFOAM 開発へ 最新の OpenFOAM をベースに開発開発の方針 特徴 GPU 実装はすべて CUDA メインループ内の処理は全て GPU で処理 GPUDirect RDMA 転送に対応現在の status 代表的なソルバを CUDA で実装本発表内容を再現することから目標高速な大規模実行の実現 GPU を用いた大規模アプリに適用しうる知見の蓄積 26
GPU での話題まとめ GPU 対応 OpenFOAM, RapidCFD の性能評価と最適化を行った 総和計算の CUDA による実装で 無駄なメモリ確保 解放コストを削減 GPU 単体では 3.3 4.4 倍高速に (CPU 比較 ) 複数 GPU でのスケーリング性能は最大で 2 倍改善 強スケーリングでの通信ボトルネック解消には GPUDirect RDMA 転送が有効 新しい GPU 版 OpenFOAM 開発へ 27
メニーコア :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 17.0.1 MPI Intel MPI 2017 update 1 コンパイルオプション -O3 -xmic-avx512 AP OpenFOAM 2.3.1(RIST 最適化版 ) 解析ソルバ pimplefoam 解析データ強スケーリング 240x130x96 cells/1-64processros 28
メモリモード変更 OpenFOAM 2.3.1 channelretau110 参考 : 京 DRAM cache MCDRAM PCG, 前処理 DIC 1 ノード内並列数の増加 (16 32 64PE) DRAM 設定時 並列数を増加させると演算処理部の時間が急増 スケーリング性能劣化 メモリモード変更 (DRAM cache/mcdram) で 3 倍弱高速に 29
クラスターモード +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
マルチノード +HyperThreading KNL/4node 構成 (quad, cache モード ) Mc: McDram HT:HyperThread xn( スレッド数 ) 参考 : 京 1ノード 2ノード 4ノード (HT:x2) (x2) (x4) (x2) 1-4 ノードまで HW モード HT に依存しないでスケール 31
KNL での話題まとめ KNL4 ノードでの OpenFOAM の性能評価を実施 メモリモード変更 (DRAM cache/mcdram)3 倍弱高速に クラスターモード HT の設定は最大で 17% 影響 Quad, SNC2, HT ON(2) が良好 4 ノードまではクラスターモード HT の設定に依存しないでスケーリングする 32
まとめ 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 (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
謝辞 本発表の内容は,HPCI システム利用研究課題 流体 粒子の大規模連成解析を用いた竜巻中飛散物による建物被害の検討 ( 課題番号 :hp170055), 課題代表者 : 菊池浩利 ( 清水建設 ( 株 ) 技術研究所 ) の高度化支援に基づくものです. 35