Transcript ポストペタスケールコンピューティングのためのヘテ

自動チューニング機構を有する
アプリケーション開発・実行環境
ppOpen-HPC
中島研吾
東京大学情報基盤センター
スーパーコンピューティング研究部門
佐藤正樹（東大・大気海洋研究所），奥田洋司（東大・人工物工学研究センター），
古村孝志（東大・情報学環/地震研），岩下武史（京大・学術情報メディアセンター），
阪口秀（海洋研究開発機構）
2011年10月11日
20111011
• ppOpen-HPCとは
• スケジュール
• H23実施状況と見通し
2
20111011
3
背 景（1/2）
• 大規模化，複雑化，多様化するハイエンド計算機環境の
能力を充分に引き出し，効率的なアプリケーションプログ
ラムを開発することは困難
• 有限要素法等の科学技術計算手法：
– プリ・ポスト処理，行列生成，線形方程式求解等の一連の共通
プロセスから構成される。
– これら共通プロセスを抽出し，ハードウェアに応じた最適化を
施したライブラリとして整備することで，アプリケーション開発者
から共通プロセスに関わるプログラミング作業，並列化も含む
チューニング作業を隠蔽できる。
– HPCミドルウェア
20111011
4
背 景（2/2）
• A.D.2000年前後
Applications
– GeoFEM，HPC-MW（奥田，中島）
– 地球シミュレータ，Flat MPI，FEM
• 現在：より多様，複雑な環境
GeoFEM, HPC-MW etc.
Fortran, C, MPI, OpenMP …
– マルチコア，GPU
– ハイブリッド並列
Vector, Scalar, MPP …
• MPIまでは何とかたどり着いたが・・・
• 「京」でも重要，全く普及しておらず
FEM code developed on PC
I/F for
I/O
I/O
– CUDA，OpenCL
– ポストペタスケールからエクサスケー
ルへ
• より一層の複雑化
I/F for
Mat.Ass.
Matrix
Assemble
I/F for
Solvers
Linear
Solver
I/F for
Vis.
Vis.
HPC-MW for T2K
I/O
Matrix
Assemble
Linear
Solver
Vis.
HPC-MW for Earth Simulator
I/O
Matrix
Assemble
Linear
Solver
Vis.
HPC-MW for
Next Generation Supercomputer
5
Hybrid並列プログラミングモデル
– OpenMP
– CUDA, OpenCL
core
core
core
memory
memory
core
core
core
core
core
core
core
core
memory
OpenMP/MPI Hybrid
core
core
core
core
core
core
core
core
memory
• Multi Threading
core
memory
– MPI
memory
• Message Passing
Flat MPI
core
core
core
core
Key-Issues for Appl’s/Algorithms
towards Post-Peta & Exa Computing
Jack Dongarra (ORNL/U. Tennessee) at SIAM/PP10
• ヘテロジニアスなアーキテクチャ：
Hybrid/Heterogeneous Architecture
– Multicore + GPU
– Multicore + Manycore (more intelligent)
– Multicore + GPU + Manycore
•
•
•
•
混合精度演算：Mixed Precision Computation
自動チューニング：Auto-Tuning/Self-Adapting
耐故障性：Fault Tolerance
通信削減アルゴリズム：Communication Reducing
Algorithms
6
20111011
7
ppOpen-HPC：概要
• ヘテロジニアスなアーキテクチャによる計算ノードを有す
るポストペタスケールシステムの処理能力を充分に引き
出す科学技術アプリケーションの効率的な開発，安定な
実行に資する
– pp= post petascale
– Five Key-Issuesを考慮
• 東大情報基盤センターに2014年度に導入予定の数10
PFLOPS級システム（ポストT2K）を最終的なターゲット：
– スパコンユーザーの円滑な移行支援
– T2K，東大次期システム（PFLOPS級），京
– TSUBAME 2.0，eScience GPUクラスタ（東大柏）
20111011
8
東大情報基盤センターのスパコン
FY
H17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
Hitachi SR11000/J2
Our Last SMP
18.8TFLOPS, 16.4TB
>50TFLOPS, >7TB
大容量メモリを使って自動並列化
MPPへ移行，センターによるサポート
Hitachi HA8000 (T2K)
140TFLOPS, 31.3TB
MPIによる並列化，メモリは遅いが通信は良い
PFLOP System w/o Accelerators
>1PFLOPS, >150TB
Hybridへの転回点，Flat MPIでも高い性能
Post T2K/Heterogeneous
>O(101)PFLOPS
本研究のターゲット
Peta
京
Exa
20111011
9
対象とするアーキテクチャ
• ホストCPU
– 8～16コア程度を有するマルチコア低消費電力CPU
• アクセラレータ
– 数十～数百のコアを有する複数のアクセラレータ
– PCI Expressに接続
• GPGPU：ホストCPUからCUDAまたはOpenCLによって利用する
• Intel Many Integrated Core（MIC）アーキテクチャに基づくメニーコア
（Knights Series）：軽量OS，コンパイラが稼働する
• GPGPU＋MICメニーコアもあり得る
• ノード間ネットワーク：Infiniband等
• アクセラレータ間通信はホストCPUを経ないで可能
• Hybrid並列プログラミングモデルは必須
20111011
10
ppOpen-HPC：基本的方針
• 対象離散化手法による分類（次ページ）
• ハードウェアに依存しない共通インタフェースを有するア
プリケーション開発用ライブラリ群，耐故障機能を含む実
行環境を提供
– 各離散化手法の特性に基づく
• 自動チューニング技術の導入により，様々な環境下にお
ける最適化ライブラリ，最適化アプリケーション自動生成
を目指す
• 数値ライブラリ，システムソフトウェア，アプリケーション
の専門家の協力：Co-Design
• 既存ソフトウェア資産（e.g. HEC-MW）の効率的利用
• 実際に動いているアプリケーションから機能を切り出す
20111011
11
対象とする離散化手法
局所的，隣接通信中心，疎行列
有限要素法
差分法
有限体積法
Finite Element Method
FEM
Finite Difference Method
FDM
Finite Volume Method
FVM
境界要素法
個別要素法
Boundary Element Method
BEM
Discrete Element Method
DEM
20111011
12
User’s Program
ppOpen-APPL
FEM
FDMii
FVM
BEM
ppOpen-MATH
MG
ii
GRAPH
VIS
MP
ppOpen-AT
STATIC
ii
DYNAMIC
ppOpen-SYS
COMM
ii
FT
ppOpen-HPC
Optimized Application with
Optimized ppOpen-APPL, ppOpen-MATH
DEM
20111011
13
ppOpen-HPC：特徴，新規性
• 汎用性・可搬性，使いやすさと高効率の両立
– 科学技術アプリの最適化→個別対応：限界がある
• 離散化手法による分類，対象の絞り込み
– それぞれの手法に応じた最適化
• データ構造の特徴を最大限利用
• 離散化手法による最適化アプローチの相違（例：疎行列）
• 自動チューニング：ppOpen-AT
– メモリアクセス最適化，Host Processor～Co-Processors負荷分散
– 新しいアーキテクチャ，計算機環境への柔軟な対応
– アプリケーション開発者だけでなくライブラリ開発者の負担も大
きく軽減できる，オープンソースアプリケーション（例：
OpenFOAM）もターゲット，他プロジェクトとの連携も容易
• 関連研究（Sphere，OpenMM，HMC）
20111011
• ppOpen-HPCとは
• スケジュール
• H23実施状況と見通し
14
20111011
ppOpen-APPL
FEM,FDM,FVM,
BEM,DEM
ppOpen-MATH
MG, VIS, GRAPH
H23年度
H24年度
H25年度
H26年度
H27年度
基本設計・プロトタイプ
整備・検証
改良
基本設計・プロトタイプ
整備・検証
改良
ppOpen-MATH
MP
基本設計・プロトタイプ
整備・検証
改良
ppOpen-SYS
COMM, FT
基本設計・プロトタイプ
整備・検証
改良
ppOpen-AT
STATIC
基本設計・プロトタイプ
整備・検証
改良
ppOpen-AT
DYNAMIC
基本設計・
プロトタイプ
整備・検証
成果物公開
国際シンポジウム
東大・ 「京」
1PFLOPS機 運用開始
運用開始
試験環境
導入
東大・
ターゲットシス
テム運用開始
改良
15
20111011
第 1 回
H24秋
第 2 回
H25秋
第 3 回
H26秋
第 4 回
H27秋
マイルストーン＝使えるソフトウェアの公開
公開内容
趣旨，目標
ターゲットマシン
ppOpen-APPL プ ロ ppOpen-APPL ， ppOpen-AT 東大T2K
トタイプ（適応格 が様々なマルチコアクラス 東大次期システム
子・動的負荷分散 タ，GPUクラスタで稼働す 京コンピュータ
を除く），領域分 る こ と を 検 証 す る 。 大 学 GPUクラスタ
割ユーティリティ，
（学部，大学院）講義への
ppOpen導入，講習会開催を継続的
MATH/VIS・
に実施する。
ppOpen-AT/STATIC
プロトタイプ
ppOpen-HPC
ppOpen-HPC の 各 機 能 東大次期システム
（ ppOpen-AT ／ （ppOpen-AT／DYNAMICを 京コンピュータ
DYNAMIC を 除 除く）をこの時期までに一 GPUクラスタ
く）
通り完成する。実アプリ
ケーションへ適用すること
により更なる改良を図る。
ポストペタスケー ターゲットとするポストペ 東大ppシステム
ス シ ス テ ム 向 け タスケールシステム向けの
ppOpen-HPCプロト ppOpen-HPCをマシン稼働開
タイプ
始と同時に利用可能とする。
ポストペタスケー 第 3 回 で 公 開 し た ppOpen- 東大ppシステム
ス シ ス テ ム 向 け HPCを実アプリケーション
ppOpen-HPC完成版 へ適用することにより，更
に改良，最適化した最終
バージョンを公開する。
16
開発環境
東大T2K
東大次期システム
GPUクラスタ（東
大）
可 能 で あ れ ば
「京コンピュー
タ」
東大次期システム
GPUクラスタ（東
大）
京コンピュータ
東大次期システム
小型メニーコアク
ラスタ
京コンピュータ
東大ppシステム
東大次期システム
小型メニーコアク
ラスタ
京コンピュータ
20111011
ターゲットマシン’s
• 東大T2K
– 平成20年6月より稼働，ピーク性能140TFLOPS
• GPUクラスタ（東大）
– 平成23年4月に東大ITCに導入されるGPUクラスタ，NVIDIA Tesla
M2050 32台搭載
• 東大次期システム
– 平成24年4月に東大ITCに導入されるスーパーコンピュータ，
ピーク性能1PFLOS級
• 小型メニーコアクラスタ
– 平成25年度前半に東大ITCに導入されるメニーコアクラスタ，本
研究の予算で導入
• 東大ppシステム
– 本研究のターゲットマシン，平成26年度に東大ITCに導入，ピー
ク性能 数十PFLOPS，pp：post-peta-scale
17
20111011
18
平成23年度
• 本研究のマイルストーンは，ターゲットとするスーパーコ
ンピュータの導入スケジュールを考慮し，前ページの表
のように設定されており，オープンソースソフトウェアの
公開を研究期間中に4回実施する予定である。
• 平成23年度は，まず平成24年秋（9月初旬を予定）の第1
回公開を目指した研究開発を実施する。
– 第１回公開では，ppOpen-APPL，ppOpen-ATが様々なマルチコ
ア，GPU環境で稼働することを検証するための最低限の機能を
満足する機能を公開。
20111011
19
第1回公開 H24秋
• ppOpen-APPLプロトタイプ（マルチコアクラスタ向け）
–
–
–
–
データ入出力部，領域間通信
簡易的な領域分割ユーティリティ
係数マトリクス生成，陽解法ソルバー
離散化手法の特性を考慮した前処理付き反復法
• ppOpen-MATH/VISプロトタイプ
– スムージングを行わない，一様差分格子を対象としたリアルタ
イム可視化処理ライブラリのプロトタイプ
• ppOpen-AT/STATICプロトタイプ
– FDM，FEM，FVM，BEM，DEMの1CPU用サンプルプログラムを
対象
– マルチコアCPU，GPU，マルチコアCPU－GPUヘテロジニアス環
境向けの自動チューニングコンパイラのプロトタイプ
20111011
20
H23実施項目：ppOpen-APPL
• 各離散化手法における必要データ，内容，格納形式等
について整理，検討
• 共通入出力インタフェースの設計
• その結果に基づき，ppOpen-APPLプロトタイプの各機能
の研究，開発，実装を行う。
• マルチコアクラスタで稼働するようなライブラリ群の研究
開発を実施
– ppOpen-MATH/VISも同様
20111011
21
H23実施項目：ppOpen-AT/STATIC
• ppOpen-APPLのサポートする各手法の1CPU用サンプル
プログラムを対象として，疎行列解法，係数マトリクス生
成部，陽解法ソルバー等に特に着目して，ソースプログ
ラムへのディレクティヴの挿入により，以下を実現するよ
うなプロトタイプの開発を実施する：
– マルチコアCPU向けに環境，問題に応じてメモリアクセス最適
化されたプログラムの自動生成
– CUDA，OpenCLによって記述され，環境，問題に応じてメモリア
クセス最適化されたGPU向けプログラムの自動生成
– マルチコアCPU－GPUヘテロジニアス環境において，環境，問
題に応じて最適な性能を発揮する資源配分決定を対象
• ppOpen-ATのためのディレクティヴについても併せて検
討を実施する。
20111011
22
H23実施項目：ppOpen-SYS
• これらとは独立に，基本的な検討，プロトタイプの開発，
実装，検証をマルチコアクラスタ，GPUクラスタを使用して
実施する。
20111011
• ppOpen-HPCとは
• スケジュール
• H23実施状況と見通し
23
20111011
24
ppOpen-APPL （1/2）
• 各手法個別の並列プログラム開発用ライブラリ群また
はテンプレート
–
–
–
–
–
①共通データ入出力インタフェース
②領域間通信
③係数マトリクス生成
④離散化手法の特性を考慮した前処理付き反復法
⑤適応格子，動的負荷分散
20111011
25
ppOpen-APPL （2/2）
• 対象分野絞り込み，各手法の特徴抽出
– FDM：非定常熱伝導，弾性動力学，自由表面付非圧縮性NS
– FVM：非定常熱伝導，圧縮性／非圧縮性NS
– FEM：定常／非定常熱伝導，弾塑性固体力学（静解析，動解
析：接触も考慮），非圧縮性NS
– BEM：地震発生サイクル，電磁場
– DEM：粉体力学，SPHにより流体も考慮，GPU化済
• ppOpen-AT/STATICとの連携
– サンプルコード提供
– メモリアクセスパターン最適化へ向けての検討
• ライブラリ化へ向けての共通データ構造，インタフェース
– netCDFベースにて検討中：専任ポスドク雇用
– H23中には基本的インタフェース決定：レガシーコードも考慮
20111011
26
ppOpen-MATH（1/2）
• ppOpen-MATH/VIS
– スムージングを行わない，一様差分格子を対象としたリアルタ
イム可視化処理ライブラリのプロトタイプをH23中に整備
• ppOpen-MATH/MG
– 計算ノード数を増やした場合のオーバーヘッド削減・・・に尽き
る
– 粗いレベルでの計算・通信手法の検討
• 通信量削減アルゴリズム
• ノード数の減少，1ノードに集める
– 現状：各MPIプロセス＝メッシュ数1になった段階で1コアに集める
27
Weak Scaling: MGCG
643cells/core, up to 512 nodes(2.05×109 cells)
T2K (Tokyo): 8,192 cores, Cray-XE6 (Hopper): 12,288 cores
T2K (Tokyo)
Cray XE6 (Hopper@LBNL)
Down is good
ICIAM11
28
Weak Scaling: MGCG
643cells/core, up to 512 nodes(2.05×109 cells)
T2K (Tokyo): 8,192 cores, Cray-XE6 (Hopper): 12,288 cores
Cray XE6 (Hopper@LBNL)
100%
100%
90%
90%
80%
80%
70%
70%
60%
60%
%
%
T2K (Tokyo)
50%
50%
40%
40%
30%
30%
CG
20%
Coarse Grid
10%
CG
20%
Coarse Grid
10%
Multigrid
0%
Multigrid
0%
64
128
256
512
1024 2048 4096 6144 8192
96
core #
192
384
768
1536 3072 6144 9216 12288
core #
ICIAM11
29
Weak Scaling: MGCG
643cells/core, up to 512 nodes(2.05×109 cells)
T2K (Tokyo): 8,192 cores, Cray-XE6 (Hopper): 12,288 cores
MPI_Isend/Irecv/Waitall，MPI_Allreduce
Cray XE6 (Hopper@LBNL)
100%
100%
90%
90%
80%
80%
70%
70%
60%
60%
%
%
T2K (Tokyo)
50%
50%
40%
40%
30%
30%
20%
20%
Comm.
10%
Comm.
10%
Comp.
Comp.
0%
0%
64
128
256
512
1024 2048 4096 6144 8192
96
core #
192
384
768
1536 3072 6144 9216 12288
core #
ICIAM11
20111011
30
ppOpen-MATH（2/2）
• ppOpen-MATH/MP
– 大気海洋カップリングライブラリ開発をまず開発し，そこから一
般的なカプラーを整備
•
•
•
•
大気： NICAM（非構造格子），MIROC（構造格子）
海洋： COCO（構造格子）
二次元（面対面）から三次元へ
本年度中に低解像度モデルを連成
– netCDFによるデータのやりとり
– H24以降
• 実行環境としての整備
– I/O，可視化，FT，M×N
• 他手法への展開
• 適切な連成アプリケーションの設定
– 津波破壊
– MHDプラズマ：格子＋粒子
提供：阪口秀博士（JAMSTEC）
正２０面体格子（NICAM）
緯度経度格子（MIROC-A）
NICAM-Agrid
NICAM-ZMgrid
HICAM-ZMgrid
MIROC-A
大気モデル1
J-cup
次世代カプラー
大気モデル2
•格子変換
•マルチアンサンブル
•IO
•Pre- and post-process
•Fault tolerance
•M×N
海洋モデル
ポストペタ計算機
システム
-システムソフト
-アーキテクチャ
COCO
領域COCO
松村モデル
Tri-Polar格子（COCO）
海洋領域モデル
非静力モデル
20111011
32
ppOpen-AT/STATIC
• 自動チューニング（AT）コンパイラ：Directive Base
–
–
–
–
最適化ライブラリ群，アプリケーションを自動生成
OpenFOAM等オープンソースアプリも対象
Hybrid，CUDA，OpenCL
メモリアクセス最適化に主眼
• ブロック化，アンローリング
– CPU～GPU負荷分散
– 基本的にはインストール時にパラメータ最適化，データベース
生成
• 各手法に適用中
– 例1：BEMにおける密行列ベクトル積（反復解法におけるボトル
ネック，Memory Bound）（ppOpen-AT/STATICの適用例）
– 例2：三次元有限要素法（予備的検討）
問題の詳細
境界要素法(Boundary Element Method,
BEM)を利用
 静電場の問題(Static Electric Filed Analysis)

◦ 地上から50cmのところに半径25cmの金属球を
配置し、金属球に＋１Vの電位を与えたとき
に金属球表面に誘起される電荷を求める
◦ 空間内の電場、ポテンシャル（電位）の値が
分かる
33
Performance Evaluation
 BEM
(Boundary Element Method)
 Main kernel is a dense matrix-vector
multiplication.
 Using BiCGStab method without
preconditioner for its iterative solver.
 Problem Sizes
◦ Small : 600 x 600
◦ Medium : 2400 x 2400
◦ Medium-Large: 15000 x 15000
◦ Large : 21600 x 21600
34
Example of The Output
Number of nodes=304 Number of faces=600
Linear system was generated.
Original relative residual norm = 1.00000000000000e+00
Step 1 relative residual norm = 5.35603994270705e-03
Step 2 relative residual norm = 5.41190843151683e-04
Step 3 relative residual norm = 2.72761060075071e-04
Step 4 relative residual norm = 5.41853759554182e-05
Step 5 relative residual norm = 5.91431206100390e-06
Step 6 relative residual norm = 3.32336208506132e-06
Step 7 relative residual norm = 9.70938119115919e-07
Step 8 relative residual norm = 7.48618319505291e-07
Step 9 relative residual norm = 5.35404930273272e-08
Step 10 relative residual norm = 1.18265378049636e-08
Step 11 relative residual norm = 8.67804885368647e-09
Relative residual norm = 8.67804896812290e-09
Time for solving liner equations : 0.149939
OK
35
ppOpen-AT Adaptation: BEM Code
A dense matrix-vector multiplication is used
in BiCGStab iteration.
 The following is main kernel with ppOpenAT description.

#pragma oat install unroll (row, col) region start
#pragma oat name MyMatVec
#pragma oat varied (row, col) from 1 to 4
for (row=0; row < dim; row++){
for (col=0; col < dim; col++){
q[row] = q[row] + mat[row][col] * p[col]; }
}
#pragma oat install unroll (row, col) region end
36
Results: BEM Code (Opteron)
T2K Open Supercomputer
 1 Core (AMD Opteron Quad core (Barcelona), 2.3GHz)
 HITACHI C compiler
 Option: -O4 –noparallel
 Whole time for part of linear equation solver

Size
Auto-tuned
Speedup
Code[Second]
Small
Original
Code
[Second]
0.0723
0.0327
2.21x
Medium
1.643
0.801
2.05x
Large
247
112
2.20x
37
Auto-tuning log : BEM Code (Opteron)
Time of The Mat-Vec Kernels
Time in second
0.0025
0.002
0.0015
0.001
0.0005
(MyMatVec
(ppOpenAT_NUMPROCS 1)
(ppOpenAT_SAMPDIST 100)
(ppOpenAT_PROBSIZE 100
(MyMatVec_I 11)
)
(ppOpenAT_PROBSIZE 200
(MyMatVec_I 16)
)
(ppOpenAT_PROBSIZE 300
(MyMatVec_I 11)
)
(ppOpenAT_PROBSIZE 400
(MyMatVec_I 11)
)
(ppOpenAT_PROBSIZE 500
(MyMatVec_I 16)
)
(ppOpenAT_PROBSIZE 600
(MyMatVec_I 16)
)
N=600
0.0022
0.0016
0.0015 0.0015
0.0014
0.0011
0.0012
0.0010
0.0011
0.0010
0.0010 0.0009
0.0009
0.0009 0.0009
0.0008
0
8 9 10 11
7 Parameter
1 2 3 4 5 6iusw:
SW12 13 14 15 16
Best: iusw=16 :
Unroll (row, col) = (4, 4)
)
38
Auto-tuning log : BEM Code (Nehalem)
(MyMatMul
(ABCLib_NUMPROCS 1)
Time of The Mat-Vec Kernels
Time in second
N=600
(ABCLib_SAMPDIST 100)
(ABCLib_PROBSIZE 100
1
(MyMatVec_I 16)
0.9
0.8599
)
0.8
(ABCLib_PROBSIZE 200
0.7099
0.7
(MyMatVec_I 16)
0.6289
)
0.6
0.5509
(ABCLib_PROBSIZE 300
0.4589
0.5
0.4469
0.4069
(MyMatVec_I 16)
0.4
0.3829
0.3759
)
0.3559
0.3429
0.3
(ABCLib_PROBSIZE 400
0.3419
0.3569
0.3639
(MyMatVec_I 16)
0.2
0.3300
0.3299 )
0.1
(ABCLib_PROBSIZE 500
0
(MyMatVec_I 16)
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
)
iusw: Parameter SW
(ABCLib_PROBSIZE 600
(MyMatVec_I 16)
)
)
Execution time
in 1000 times
Best: iusw=16 :
Unroll (row, col) = (4, 4)
39
ppOpen-AT Adaptation: BEM Code



CPU version
OpenMP Parallelization + Unrolling
One Dimensional Implementation
#pragma oat install unroll (row, col) region start
#pragma oat name MyMatVec
#pragma oat varied (row, col) from 1 to 4
#pragma omp parallel for private(col)
for (row=0; row < dim; row++){
tmp = 0.0;
for (col=0; col < dim; col++){
tmp = tmp + mat[row*dim+col] * p[col];
}
q[row] = tmp;
}
#pragma oat install unroll (row, col) region end
40
ppOpen-AT Adaptation: BEM Code

CPU + GPU kernel selection
#pragma oat select region start
#pragma oat name SelectGPU
#pragma oat select sub region start
pbicgstab(dim, a, rhs, sol, tor, max_steps);
#pragma oat select sub region end
#pragma oat select sub region start
pbicgstab_gpu(dim,a, rhs, sol, tor, max_steps);
#pragma oat select sub region end
#pragma oat select region end
41
BEM : GPU code with PGI accelerator
#pragma acc data region copyin (mat[0:dim-1][0:dim-1], shdw[dim] ) {
/* BiCGSTAB iteration */ for (step=1;step<=max_steps; step++) {
#pragma acc region {
for (row=0; row < dim; row++){ ap[row] = 0.0; }
for (row=0; row < dim; row++){
for (col=0; col < dim; col++){
ap[row] = ap[row] + mat[row][col] * p[col]; } }
for(i=0;i<dim;i++){p[i] = r[i] + beta * ( p[i] - zeta * ap[i] );}
/* No preconditioning */ for (i=0;i<dim;i++){kp[i]=p[i];}
for (row=0; row < dim; row++){ akp[row] = 0.0; }
for (row=0; row < dim; row++){
for (col=0; col < dim; col++){
akp[row] = akp[row] + mat[row][col] * kp[col]; } }
} //************ end of acc region
nom = dot_product(dim, shdw, r); den = dot_product(dim, shdw, akp); alpha = nom / den; nomold = nom;
#pragma acc region {
for(i=0;i<dim;i++){t[i] = r[i] - alpha * akp[i];}
/* No preconditioning */ for (i=0;i<dim;i++){kt[i]=t[i];}
for (row=0; row < dim; row++){ akt[row] = 0.0; }
for (row=0; row < dim; row++){
for (col=0; col < dim; col++){
akt[row] = akt[row] + mat[row][col] * kt[col]; } }
} //************ end of acc region
nom = dot_product(dim, akt, t); den = dot_product(dim, akt, akt); zeta = nom / den;
for(i=0;i<dim;i++){sol[i] = sol[i] + alpha * kp[i] + zeta * kt[i];}
for(i=0;i<dim;i++){r[i] = t[i] - zeta * akt[i];}
beta = alpha / zeta * dot_product(dim, shdw, r) / nomold;
rnorm = sqrt( dot_product(dim, r, r) );
if (rnorm/bnorm < tor){break;}
} //********* end of data acc region
} /*********************************/
/* Confirmation of residual */ residual_direct(dim, mat, sol, rhs, r); rnorm = sqrt( dot_product(dim, r, r) );
}
42
Results: BEM Code (Nehalem+C2050)


Nehalem + C2050
CPU: OpenMP Parallel Threads Execution for CPU
◦ #Threads: 8, One dimensional Optimization for Mat-Vec


GPU: Optimized with PGI Accelerator
PGI Compiler
◦ -Msafeptr=all -fastsse -lm -mp -ta=nvidia -Mlist -Minfo=accel,mp

Whole time for part of linear equation solver
Size
CPU
Code
[Second]
GPU Code
[Second]
Speedup
Small
0.002
0.022
0.09x
Medium
0.128
0.125
1.02x
Medium-Large 7.96
6.99
1.13x
Large
Out of
Memory
-
-
43
Example of FEM


FEM code for 3D Solid Mechanics
My program is based on existing FEM application
for CPU



Optimization is applied for CPU



Z
3x3 blocked
Diagonal/Lower/Upper data are
stored in each individual array
Force
CPU:Nehalem 2.67GHz


based on GeoFEM (http://geofem.tokyo.rist.or.jp/)
C language, CRS(CSR) format
42.7GF, 32GB/s
GPU: Tesla C2050

515GF, 144GB/s
Y
X
Target for Optimization
• Sparse Linear Solver (Preconditioned Iterative
Solver)
• Matrix Assembling
201109
45
Performance of SpMV for FEM
Computation
[Ohshima et al., 2011]
• GeoFEM Benchmark
• CPU: Nehalem 2.67GHz (quad-core)
Z
– 42.7GFLOPS, 32GB/s
– 3.5 GFLOPS with 4xOpenMP Threads
• GPU: Tesla C2050
– 515GFLOPS, 144GB/s
– 11.4 GFLOPS with 240 Threads
• Optimized implementation based on cusp
X
Y
• GPU is 3.25x is faster than CPU
• (GPU+CPU) is 30% faster than GPU only !!
201109
46
Parallel Matrix Assembling by Coloring

Parallel calculation with coloring: multi-color
method
1.
2.
mark the elements which can be calculated at same
time as "same color" (coloring)
parallel calculation in each color (calculation)
47
Matrix Assembling
106 elements, 8 colors
• CPU: 4.73 sec.
• GPU: 1.28 sec. (x 3.46)
201109
48
20111011
49
疎行列ソルバーによるアプリケーション
• CPU～GPUの性能差が現状ではそれほど大きくない（1ソ
ケット）
– 両者を利用することが効率的
– 適切な負荷分散，CPU～GPU間の通信削減
• Multicore～Manycore（e.g. MIC）
• 差分法
– 構造格子，データパターンが予測可能
• 有限要素法，有限体積法
– 非構造格子
– データによって最適な解が異なる可能性
– 実行時の試行はある程度必要となる可能性：最小限にしたい
20111011
50
その他
• ppOpen-MATH/GRAPH
– スケーラブルなグラフ分割アルゴリズム
• ppOpen-SYS/COMM
–
–
–
–
ppOpen-HPCで使われるアプリに限定した通信ライブラリ
MPIと同じインタフェース
ホストCPUのメモリを経由しない通信
MIC内のプログラミングモデル
• スレッドonly，スレッド＋メッセージパッシング
• ppOpen-SYS/FT
– リスタートファイルの効率的IO
20111011
51
状 況
• H24-9月の公開に向けて順調な進捗
• 非構造格子向けのppOpen-ATの戦略がカギ
• H23年度中に予定していた国際ワークショップは延期
– 2012年10月以降
– 若手中心
• 企画，運営
• 招待講演者
– 各チームからもご協力（講演等）をお願いしたい