科 学 技 術 計 算 分 科 会

選 出

科学技術計算分科会 2012 年度会合 より

HPC 分野における GPU 活用技術
−アクセラレータ技術 WG 成果報告−

井上 弘士（九州大学）

2012 年 10 月 24 日

サイエンティフィック・システム研究会
科学技術計算分科会 2012 年度会合

HPC 分野における GPU 活用技術
〜アクセラレータ技術 WG 成果報告〜
SS 研 アクセラレータ技術 WG
とりまとめ役：井上弘士（九州大学）[アブストラクト]
本講演では、2009 年からおよそ 2 年に渡り実施したアクセラレータ技術 WG の成果を報告する。特に、アクセラレーション技術として注目を集めている GPU の活用に焦点をあて、
本 WG における主な成果の 1 つである「今後の GPU 活用のあるべき姿」を探るために実施した評価実験について報告する。
[キーワード]アクセラレーション、GPU、ハイパフォーマンスコンピューティング、コンパイラ、チューニング
１．はじめにスパコンシステムの計算性能は、これまで概ね 3 年で 10 倍のペースで向上してきた。近年はプロセッサの動作周波数を向上させることが技術的に難しくなったため、プロセッサコアの単体性能を向上させるのではなく、スパコンを構成するノード数を増やす、プロセッサのコア数を増やす（マルチコア化）など、数を増やすアプローチで性能向上のペースを維持してきた。その中で、近年注目されているのがアクセラレータである。汎用プロセッサは、HPC 以外の様々な用途での使用を想定しており、様々なプログ
ラムで簡単に安定した性能が得られる様に実装面積の半分以上をキャッシュメモリとして使用しており、演算ユニットの占める面積比率は低い。それに対して、アクセラレ
ータは HPC 用途にフォーカスしたプロセッサで、キャッシュメモリなどの非演算ユニットのリソース量を最小限に抑え、多数の演算ユニットを搭載して、演算ユニットの面積比率を高めることで、演算性能を高めるアプローチのプロセッサである。アクセラレ
ータは高い演算性能を持つ反面、プログラム開発が難しいなど、その性能を引き出すのは簡単でない。実際に、特定アプリでは高い性能が得られるが、それ以外のアプリでは高い性能が得られないという問題がある。しかし、エクサスケール世代では飛躍的に電力あたり演算性能を高める必要があるため、汎用プロセッサと比べて電力あたり演算性能の優れたアクセラレータ技術は、エクサスケール世代の必須技術であると期待されている。
1

このような背景の下、SS 研アクセラレータ技術 WG を立ち上げ、アクセラレータとして近年普及しつつある GPU を事例として、将来のアクセラレータのあるべき姿を模索・議論した。特に、GPU に代表されるアクセラレータ上でのプログラム開発の現状とその課題、また、どうすればアクセラレータで高い性能を得られるのか、具体的な事例を用いて調査した結果を報告する。
２．WG 活動概要
WG の活動を支えたメンバーを以下に示す。活動期間中に 8 回の会議を実施した。これ
らの会議においては、各メンバーによるアクセラレーション技術に関する技術報告に加え、各種実験トライアルの結果報告と議論、さらには、今後のアクセラレーション技術（大規模センター等での運用技術も含む）に関する議論を展開した。

氏名

所属

担当幹事 村上 和彰

九州大学推進委員 青木 尊之

東京工業大学

遠藤 敏夫

東京工業大学黒川 原佳

理化学研究所

滝沢 寛之

東北大学

伊野 文彦

大阪大学

井上 弘士

九州大学

本田 宏明

九州大学

堀田 耕一郎富士通(株) 次世代テクニカルコンピューティング開発本部

丸山 拓巳

富士通(株) エンタプライズサーバ事業本部

坂口 吉生

富士通(株) TC ソリューション事業本部

佐々木 啓

富士通(株) TC ソリューション事業本部

久門 耕一(株)富士通研究所

成瀬 彰(株)富士通研究所

３．おわりに

本 WG の活動において、様々な知見を得ることができた。特に、1)市販コンパイラを用いた「手軽な」GPU の活用においては、ディレクティブの挿入に 2〜3 時間程度費やすだけで、人手でのチューニングと比較して 40%程度の性能を得た。また、2)一般には「GPU に不向き」とされるアプリケーションにおいても、実装アルゴリズムを GPU向けに再検討することで大幅な性能向上を達成できた。今後は、アプリ開発者と実装チューニングの連携、GPU−CPU の協調実行、などが重要になる事が分かった。

2

発表内容
HPC分野におけるGPU活用技術〜アクセラレータ技術WG成果報告〜
SS研 アクセラレータ技術WG取りまとめ役 井上弘士（九州大学）• ワーキンググループ（WG）設置の背景• WG活動概要• GPUを「手軽に」使って性能が改善するか?
• GPUに不向きと⾔われるアプリは本当に不向きなのか?
• 将来のアクセラレータ活用のあるべき姿
は?

1

2

エクサ実現へのチャレンジ• エクサ・フロップス級スパコンの実現における最大の課題は?
→「消費電⼒の壁」• 1EFlops@20MWスパコンを実現するには?
→「京」を基準にすると消費電⼒2xで性能100x
1E+10
1E+09

Gflops (Rmax)

WG設置の背景

1EFlops

100x

1E+08
1E+07

Top100@2011

2x

1E+06
1E+05
1E+04

3

DARPA IPTO Exascale Computing Study
http://www.darpa.mil/tcto/docs/ExaScale_Study_Initial.pdf

説明

システム例

– ITRS@2006

高性能プロセッサ搭載
型ノード

Red Storm

Intel, AMDなどの
プロセッサ

Light Node (LN)

組込み向けプロセッサ搭載型ノードBlue Gene/L, /P

PPC440, 450

Aggressive Node (AN)1EFlopsのシステムを前提とした仮想ノード

無し

無し

10000

20MW

4

• ノード消費電⼒モデル– Simplisticlly Scaled Power Model• プロセッサ→ITRS@2006• 主記憶（DRAM）→∝#DRAMチップ• ルータ→ベースと同じ 割合で一定• メモリ総容量– ∝理論ピーク性能– Fully Scaled Power Model• プロセッサ→ITRS@2006• 主記憶（DRAM）→∝#DRAMチップ×性能• ルータ→∝性能

プロセッサ例

Heavy Node (HN)

1000

Power (KW)DARPA IPTO Exascale Computing Study• 半導体デバイス• 既存マシンをベースにしたスケーリングによりシステム性能を予測• 消費電⼒と⾯積（最⼤ラック数）制約を考慮
ノード名

100

コア

ラック

メモリ(DRAM) ルータ

Cache

5

プロセッサ

ノード

ボード

システム

6

本WGの狙いDARPA IPTO Exascale Computing Study• 単純な量的変化だけではエクサ@2018は到達できない!

• 現在最も有望なアクセラレーション技術→GPU– Fullyで50〜60x，Simplisticallyで10xの性能向上が必要• 何らかの質的な変化が必要では?

– 「GPU向きアプリ」を対象とした多くの事例報告– アクセラレーション技術

Figure 7.14: Light node strawman performance projections

1EFlops50〜60x

消費電力制約無し∞W Simplistically20MW Simplistically∞W Fully
20MW Fully

10x– GPUを「手軽に」使って性能が改善するか?

• PGIコンパイラを用いた「ライトな」チューニング– GPUに不向きと⾔われるアプリは本当に不向きなのか?
• GPUに不向きとして知られるアプリを対象に「ヘビーな」チューニング• 上記の調査を通して将来のアクセラレータ活用のあるべき姿を模索・議論

LN: Light Nodeの結果
2018• 素朴な，でも，本質的な疑問

87

メンバーと活動内容
氏名

WG活動概要

所属

担当幹事

村上 和彰

九州大学

推進委員

⻘⽊ 尊之

東京工業大学

遠藤 敏夫

東京工業大学⿊川 原佳

理化学研究所

滝沢 寛之

東北⼤学

伊野 文彦

大阪大学

井上 弘士

九州大学

本田 宏明

九州大学

堀田 耕⼀郎富士通(株)丸山 拓⺒

富士通(株)

坂口 吉生

富士通(株)佐々木 啓

富士通(株)

久門 耕一(株)富士通研究所

成瀬 彰(株)富士通研究所

会合

日時

第一回

2009年
10月29日(木) WG全体スケジュールの議論
 今後の議題検討

活動内容

第二回

2010年
1月19日(火) 会員報告 (PGIアクセラレータコンパイラ、コーンビー
ム再構成の高速化、初心者によるGPUプログラミング
事例、複合型計算機向けソフトウェア開発環境)
 富士通報告 (GPGPUプログラミングの状況、GPUSIMから⾒たGPU)

第三回

2010年
4月7日(水) 会員報告 (Linpack on GPU搭載スパコンTsubame) 富士通報告 (OpenCLの評価)
 OpenCLに関する議論

第四回

2010年
7月14日(水) 会員報告 (PGIコンパイラ評価)
 富士通報告 (nVidia新GPU評価)

第五回

2010年
12月3日(⾦)

 会員報告 (アクセラレータの大規模システム導入課題、Linpack on Tsubame2)

第六回

2011年
3月25日(⾦)

 会員報告 (PGIコンパイラ評価)
 富士通報告 (分子軌道法プログラムのGPU化検討)

第七回

2011年
8月19日(⾦)

 会員報告 (CUDA Fortran評価)
 富士通報告 (分子軌道法プログラムのGPU化結果)

第八回

2012年
1月19日(木)

 まとめ

9

10

3人の開発者による
PGIアクセラレータ利⽤事例• PGIアクセラレータ：– nVidia GPU向けのディレクティブ挿入方式のプログラム
開発言語・環境• 3人の開発者

GPUを「手軽に」使って性能が改善するか?

開発者A

OpenMP・MPIによる並列プログラム開発経験はあるが，PGIアクセ
ラレータの使用経験は無し．

開発者B

計算科学分野のアプリ開発経験は豊富．逐次プログラム開発が主体であり，プログラムの並列化はあまり詳しくない．

開発者C

並列処理のエキスパートであり，CUDA・OpenCLでのプログラム開発経験がある．

• 3種類のアプリ

11– 姫野BMT– 2D‑FDTD– 分子軌道法

12

実験結果
姫野BMT
GFLOPS

Directive数

PGI ver.1

14.9

2⾏

PGI ver.2

19.2

CUDA版
(by 開発者Ｃ)

2D‑FDTD
開発工数

分子軌道法

GFLOPS

Directive数

1時間

18.5

10⾏

8⾏

+30分

…

…

…

…

50

…

数日

…

…

…

…

CUDA版
(by 開発者B)

…

…

…

45.5

…

3⽇程度

…

Xeon X5670
2.93GHz 1コア

3.9

…

…

7.0

…

…

…

Xeon X5670
2.93GHz 4コア

7.8

…

…

…

…

…

…

(学習に数日)

開発工数

GFLOPSGPUコード
生成できずGPUに不向きと⾔われるアプリ
は本当に不向きなのか?

エキスパート（開発者C）のCUDA版実装と比較して40%程度の性能
1314

GPUでの高速化は難しい「らしい」

分子軌道法とは分子軌道法プログラムの GPU 化は一般的に難しいとされている• 数例の報告のみ• 分子内において，主に電子がどのような運動をしており，どのようなエネルギーを持っているかを量⼦化学的計算により求める方法の一つ• オリジナルコード– コードが複雑、ステップ数大– PGI アクセラレータコンパイラのディレクティブの方法では、GPU へのオフロードコードは生成されず– 分子物性の解析– 創薬、新素材の開発– ex) プリンタのカラーインク、液晶ディスプレイ• GPU使用の問題点– 倍精度計算の使⽤– 開平逆数、指数関数計算が必要– 条件分岐計算– レジスタ量が不⼗分– シェアードメモリを利⽤のための適切な⽅法が不明確– デバイスメモリサイズの制限– などなど16

15

ボトルネックはG⾏列計算

分⼦軌道法プログラムの処理フロー
分子の座標、座標、
電子数、基底関数
データなどを入力
重なり積分S, １電子ハミルトン行列
Hcore、および、２電子積分を計算
重なり積分をCholesky分解（S=UTU）
密度行列の初期値（D0）を設定

CPU• 積分駆動型アルゴリズムによる２電⼦フォック⾏列計算• Obara のアルゴリズムによる積分計算法計算• 全てC言語で記述

Core-i7 2600K
3.4GHz (1core)Compiler
Intel v12.1 (-O3)

Program

密度行列と２電子積分から２電子
ハミルトン行列Gを計算
密度行列を更新

Kyudai-HFR

Input data

F=Hcore+Gを計算

h16o8 (24 atoms)

FC=SCεを解く

Elapsed time

Cから新たな密度行列Dを計算

no

密度行列が収束した？
yes
収束したF, C, Dなどを用いて、望みの期待値（エネルギーなど）を計算

分子軌道法プログラムの高速化≒ G⾏列計算の⾼速化

2電子積分計算（ERI）を含む
G行列計算式1714.29 sec

18

G⾏列計算の概要＆OpenMPによるCPU向け並列化

係数⾏列作成1st loop begin
2 nd loop begin

3 rd loop begin

omp parallel for

1st loopのiterationをスレッドに分配4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算

CPU向け並列化

G⾏列更新
2 nd loop end
1st loop end

omp atomic

他スレッドと同じ箇所を更新する可能性
G [ ij ] += 4 * x * D [ kl ]G [ kl ] += 4 * x * D [ ij ]G [ ik ] ‑=x * D [ jl ]G [ il ] ‑=x * D [ jk ]G [ jk ] ‑=x * D [ il ]G [ jl ] ‑=x * D [ ik ]19

20

OpenMPによる並列化効果

実際には6種類の積分計算(コア数:1〜4)
CPU
Core‑i7 2600K
3.4GHz (1‑4 cores)(ss,ss)タイプ

Compiler(ps,ss)タイプ

Intel v12.1(‑O3 ‑openmp)(ps,ps)タイプ

Program(pp,ss)タイプ(pp,ps)タイプ(pp,pp)タイプ

Kyudai‑HFRCPU: Core‑i7 2600K 3.4GHz (1core)Compiler: Intel v12.1 (‑O3)

6種類の積分型を全てOpenMPで並列化

Input datah16o8 (24 atoms)

Program: Kyudai‑HFR
Input data: (H2O)8 (24 atoms)
Elapsed time: 14.29 sec

21

22

性能

実⾏時間（1コアに対する性能向上）
積分型(ss,ss)(ps,ss)(ps,ps)(pp,ss)(pp,ps)(pp,pp)

CPU

GPU

1コア

4コア

246
(1.0x)
279(1.0x)
122(1.0x)
67(1.0x)
62(1.0x)
12(1.0x)

66.2(3.7x)
85.8(3.3x)
39.6(3.1x)
21.3(3.1x)
21.0(3.0x)
4.0(3.0x)?

GPU向け並列化（NAÏVE実装）

2324

G⾏列計算処理 ((ss,ss)積分型)
係数⾏列作成1st loop begin
2 nd loop begin
3 rd loop begin
4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算G⾏列更新
2 nd loop end
1st loop end

G⾏列計算のGPGPU化

係数⾏列作成係数⾏列数：16種
総量：2.4MB (⼊⼒依存)
1st & 2 nd loop：
ループ回数：139,656 (⼊⼒依存)
ERI計算：
ループ回数：81 (⼊⼒依存)
ERI値：1種 (積分型依存)
G配列更新
前出ERI値と関連する部分を更新
G⾏列更新：6箇所 (積分型依存)
G配列要素数：1,596 (⼊⼒依存)(*) Input data: (H2O)8 →24 atoms

係数⾏列作成1st loop begin
2 nd loop begin
3 rd loop begin
4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算(*) Input data: h16o8 (24 atoms)

ERI計算：
ループ回数：81 (⼊⼒依存)
ERI値：1種 (積分型依存) マルチスレッド化
 各ERI値の計算には5スレッド

使⽤（定量的な解析より）

3 rd loop begin
4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算

G⾏列更新

G⾏列更新
2 nd loop end(*) Input data: h16o8 (24 atoms)

1st loop end

G⾏列計算のGPGPU化

2 nd loop end
1st loop end

Nスレッド並列(*) N=2〜81

N‑1スレッドは待機（6回のatomic操作）(*) Input data: h16o8 (24 atoms)

27

28

GPU環境

係数⾏列作成1st loop begin
2 nd loop begin

G⾏列更新

26

係数⾏列作成1st loop begin
2 nd loop begin

2 nd loop end

3 rd loop begin
4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算

1スレッドに割当

ボトルネックはどこに?

G⾏列計算のGPGPU化
係数⾏列作成1st loop begin
2 nd loop begin

1st loop end

 1st & 2 nd loopの各iterationを

G⾏列更新
2 nd loop end
1st loop end

25

3 rd loop begin
4 th loop begin
初期積分計算垂直漸化式計算4 th loop end
3 rd loop end
水平漸化式計算

1st & 2 nd loop：
ループ回数：139,656 (⼊⼒依存)(*) CC = CUDA Core

 Tesla C2050

Stream Multi-processor
Register Files (128KB)

Fermi世代
部分G⾏列更新
ERI値毎に、関連する箇所を更新
G⾏列更新箇所：6(全要素数：1,596)

CC
CC
CC
CC

 CUDA 4.2(x14)

L2(128KB)

 atomicCAS()による排他処理CAS: Compare And Swap

MemoryController(*) Input data: h16o8 (24 atoms)

GDDR5

29

CC
CC
CC
CC

CC
CC
CC
CC

CC
CC
CC
CC

CC
CC
CC
CC

CC
CC
CC
CC

SM

X-Bar

 他スレッドが同時に同じG⾏列要素を更新する可能性

CC
CC
CC
CC

L1+Shmem:64KB
SM 更新箇所は不連続

CC
CC
CC
CC

64bit

(x6)

L2(128KB)
MemoryController
64bit
GDDR5

30

性能

実⾏時間（1コアに対する性能向上）
積分型(ss,ss)(ps,ss)(ps,ps)(pp,ss)(pp,ps)(pp,pp)

CPU

GPU

1コア

4コア

246
(1.0x)
279(1.0x)
122(1.0x)
67(1.0x)
62(1.0x)
12(1.0x)

66.2(3.7x)
85.8(3.2x)
39.6(3.0x)
21.3(3.1x)
21.0(2.9x)
4.0(3.0x)

Naïve

45(5.4x)
‑‑‑

?

‑‑‑
‑‑‑

GPU向け並列化（実装アルゴリズムの最適化）

‑‑‑
‑‑‑
31

ボトルネックはどこに存在する?
〜atomic操作を外すと…〜

32

Kernel分離：atomic操作を無くす！
1st & 2 nd loop begin• CPU (1core)
246 ms• GPU (2x2x5)
45 ms (5.4x)• GPU (w/o atomic操作) 25 ms (9.8x)

1st & 2 nd loop begin

ERI計算

ERI計算
(ERI値をメモリにWRITE)

分離

G⾏列更新

1st & 2 nd loop end

1st & 2 nd loop end+

ERI値の計算直後に関連する
G⾏列要素を更新(*)ただし、正しい結果は得られない

1st & 2 nd loop begin(ERI値をメモリからREAD)

G⾏列更新

ボトルネックはG⾏列計算のatomic操作→これをなんとかしないと速くならない!

1st & 2 nd loop end

ERI値をメモリに記録するオーバーヘッドは高いメモリバンド幅で隠蔽(と期待)
33

34

性能

性能

実⾏時間（1コアに対する性能向上）
積分型(ss,ss)(ps,ss)(ps,ps)(pp,ss)(pp,ps)(pp,pp)CPU

実⾏時間（1コアに対する性能向上）

GPU

1コア

4コア

246
(1.0x)
279(1.0x)
122(1.0x)
67(1.0x)
62(1.0x)
12(1.0x)

66.2(3.7x)
85.8(3.2x)
39.6(3.0x)
21.3(3.1x)
21.0(2.9x)
4.0(3.0x)

Naïve

積分型

カーネル分割45(5.4x)
‑‑‑
‑‑‑
‑‑‑
‑‑‑
‑‑‑

10.5
(23.4x)16.2(17.2x)12.6(9.7x)
6.3(10.6x)10.1(6.1x)
9.6(1.3x)

CPU
1コア

35

Naïve

246
66.2
45
(1.0x)(3.7x)(5.4x)
279
85.8‑‑‑(ps,ss)(1.0x)(3.2x)
122
39.6‑‑‑(ps,ps)(1.0x)(3.0x)
67
21.3‑‑‑(pp,ss)(1.0x)(3.1x)• 積分型により性能UP率に⼤きな違い
62
21.0
‑‑‑
(pp,ps)• 複雑な積分型はGPUで性能向上難しい?
(1.0x)(2.9x)→さらなるチューニングを試みる!
12
4.0
‑‑‑
(pp,pp)(1.0x)(3.0x)(ss,ss)

?

GPU
4コア

カーネル分割

10.5 GPU向き(23.4x) •繰り返し多
•計算単純
16.2
(17.2x)12.6(9.7x)
6.3(10.6x)10.1
GPU不向き(6.1x) •繰り返し少9.6 •計算複雑(1.3x)
36

ERI計算の複雑さは積分型次第

積分型によるG⾏列計算処理の違い
係数行列作成1st loop begin
2nd loop begin

積分型

3rd loop begin
4th loop begin
初期積分計算垂直漸化式計算
4th loop end
3rd loop end
水平漸化式計算G行列更新
2nd loop end
1st loop end

計算されるERI値(ss,ss)
1個(pp,pp) 81個更新されるG配列要素(ss,ss) 最多 6箇所(pp,pp) 最多486箇所(ss,ss)(ps,ss)(ps,ps)(pp,ss)(pp,ps)(pp,pp)# usedregisters
54
63
63
63
63
63

Shmem Spill stores Spill loads(bytes)(bytes)(bytes)
0
0
0
0
64
64
0
296
220
0
164
136
0
852
636
0
2,708
2,332

• レジスタ数が足りない→レジスタスピル発生

37

38

性能

Shmemの活用

実⾏時間（1コアに対する性能向上）

積分型(ss,ss)(ps,ss)(ps,ps)(pp,ss)(pp,ps)(pp,pp)# usedregisters
50
63
62
63
63
63

Shmem Spill stores Spill loads(bytes)(bytes)(bytes)
0
0
0
0
28
28
4,224
164
116
768
224
44
7,680
700
448
7,872
232
44

39

(ps,ss)(ps,ps)(pp,ss)(pp,pp)

CPU

GPU

1コア

4コア

246
(1.0x)
279(1.0x)
122(1.0x)
67(1.0x)
62(1.0x)
12(1.0x)

66.2(3.7x)
85.8(3.2x)
39.6(3.0x)
21.3(3.1x)
21.0(2.9x)
4.0(3.0x)

Naïve

カーネル分割45(5.4x)
‑‑‑
‑‑‑
‑‑‑
‑‑‑
‑‑‑

10.5
(23.4x)16.2(17.2x)12.6(9.7x)
6.3(10.6x)10.1(6.1x)
9.6(1.3x)

更なるチューニング

10.5(23.4x)14.9(18.8x)8.6(14.1x)4.2(15.9x)9.2(6.8x)
7.6(1.5x)
40

各種⼊⼒データでの性能比較（vs. 1コア）

全体性能比較(CPU vs GPU)
14.3秒

(ss,ss)(pp,ps)• Shmemを使ってレジスタスピル量を削減– 他にも細かなチューニング

積分型

CPU
Core-i7 2600K
3.4GHz (1core)Intel compiler v12.1
(-O3)

GPU
Tesla C2050
CUDA 4.2

1.16秒

ProgramKyudai-HFR

CPU(1core)GPU: 12.3x

Input data: h16o8(24 atoms)
41

42

各種⼊⼒データでの性能比較（vs. 4コア）

将来のアクセラレータ活用のあるべき姿は?
44

43

GPUを「手軽に」使って性能が改善するか?

GPUに不向きと⾔われるアプリは本当に不向きなのか?

• GPUに向いている単純なコード• 実装アルゴリズムをGPU特性に合わせて⾒直すことでCPUを大きく上回る性能を達成– 数⾏のディレクティブ挿⼊– エキスパートによるチューニング実装と比較
して40%程度の性能– GTX 580性能：対CPU(1core)で最大21.2倍
対CPU(4core)で最大6.0倍• コンパイルできないコードも存在• 本トライアルで性能向上を達成できた理由– 理解し易いシンプルなコードの提供– 実装屋の存在

45

CPU

CPU vs. GPU
Memory(on-chip)Memory(on-chip)アクセラレータ「活用の」あるべき姿（1/2）

GPU
如何に上手
く使うか?
如何に稼働率を高めるか?

Control

ControlCalculation

Calculation• GPU化に「向く」部分と「向かない」部分はある– 依然としてアプリ屋と実装屋のギャップは大きい• アクセラレータ活用のあるべき姿– 素性の良いアプリ→コンパイラが「そこそこ」使える– 向かないアプリ→アプリ屋と実装屋の協調（Co‑
Design?）で多くを「向く」アプリに!

如何に性能向上に活かすか?

Memory46

理解し易い
シンプルなコード（マシン非依存）

Memoryアプリ・レベルの最適化に注⼒47

アーキ依存な
⾼度最適化

実装レベルの最適化に注⼒
フィードバック

アクセラレータHWとしての進化
性能解析

48

アクセラレータ「活用の」あるべき姿（2/2）• CPU(Fat Core)とGPU(Thin Cores)の併用– どうしてもGPUに向かない処理はある（pppp型など）– ただし，データサイズが⼩さく実⾏時間が短いことが多い– このような処理はCPUに任せる• タスク並列的なCPU/GPU併用

49

科 学 技 術 計 算 分 科 会

選 出

科学技術計算分科会 2012 年度会合 より

スーパーコンピュータ「京」でのMPI の実装と評価

三浦 健一（富士通株式会社）

サイエンティフィック・システム研究会
科学技術計算分科会 2012 年度会合

2012 年 10 月 24 日

スーパーコンピュータ「京」でのMPI の実装と評価
三浦 健一
富士通株式会社 次世代テクニカルコンピューティング開発本部 ソフトウェア開発統括部[アブストラクト]
スーパーコンピュータ「京」1の共用が 9/28 から開始された。
「京」で使用されている MPI
ライブラリでは 8 万ノード規模の並列計算に耐えられ、かつ高速通信ができるよう、様々な工夫が施されている。本発表においては、「京」における MPI ライブラリの特徴である、
省メモリ通信、通信最適化、最適ランク配置、低レベル通信ライブラリの実装を説明し評価する。また併せて、アプリケーションへの適用事例に関して報告する。
[キーワード]Tofu インターコネクト，MPI 通信，大規模並列，省メモリ，最適ランク配置

1

はじめに
スーパーコンピュータ「京」の共用が 9/28 から開始された。「京」は 8 万ノード超のノード群が Tofu(Torus fusion)と呼ばれる 6 次元メッシュ/トーラス型インターコネクトで結合された構成となっている。Tofu では 6 次元メッシュ/トーラスを生かした故障ノード回避や迂回通信、ジョブ単位に論理トーラス割当を行うことが可能となり高い運用性を実現している。また Tofu では 4 個の通信インターフェイスを用いた同時通信や高機能バリアを用いたハードウェア同期を行うことでより高速に通信を行うことが可能である。我々は「京」上で Tofu を用いて通信を行う MPI ライブラリの開発を行い、大規模環境での評価を行った。

2

MPI ライブラリの実装概要「京」の MPI ライブラリは Open MPI をベースに拡張を行ったが、8 万ノード規模の Tofu

インターコネクトへの適用にあたり、以下の点に注意して設計を行った。


8 万ノード規模の並列化に耐えうる実装

Tofu の特性を生かした高速通信

1「京」は理化学研究所の登録商標です。



ユーザの使いやすさ
MPI 通信を行うためには通信相手毎に送受信バッファが必要となるが、8 万ノード分の送受信バッファを確保していたのでは MPI ライブラリ自体のメモリ使用量が膨大になってしまう。
「京」の MPI ライブラリでは省メモリ通信モードと高速通信モードを使い分けること
で、省メモリを維持しつつ高速に通信を行うことを狙っている。
Tofu の特性を生かした高速通信では、特に大規模並列計算で問題となる集団通信に関して、Tofu 向けアルゴリズムを新規開発することで、ハードウェア性能を最大限引出す事を狙っている。
またユーザの使いやすさとして、通信を行うランクを近くに配置することで性能チュー
ニングを簡単に行うことができる RMATT(Rank Map Automatic Tuning Tool)、突き放し通信
や 4 個の通信インターフェイスを使った並列通信を容易に記述することができる「拡張 RDMA
インターフェイス」の設計・実装を行った。

3

MPI ライブラリの評価
省メモリ通信では 8 万ノード規模の並列処理においてデフォルトメモリ使用量を 400MB

程度まで抑えることができた。通信最適化においては集団通信の Tofu 専用アルゴリズムの採用により既存アルゴリズムに比べ 5〜10 倍程度の性能向上が確認できた。RMATT を用いた
最適ランク配置では、NAS parallel Benchmark に適用することで 7%の性能改善が確認できた。「拡張 RDMA インターフェイス」では 4 個の通信インターフェイスを用いた通信を行うことで性能改善の効果が確認できた。
アプリケーションへの適用事例としては昨年度ゴードン・ベル賞を獲得した RSDFT や 4
部門全てで 1 位を獲得した HPCC の Global FFT や Global RandomAccess の通信ライブラリとして使用され高効率の実行に貢献した。

4

まとめ「京」でサポートしている MPI ライブラリの実装と評価に関して述べた。
「京」の運用は始まったばかりであり、今後多くの技術者が「京」に触れて「京」の素晴らしさを体感していただけることを期待している。本 MPI ライブラリを用いることで大規模アプリケーシ
ョンの開発・運用に少しでも貢献できれば幸いである。
[参考文献](1) 住元真司 et al., 「京」のための MPI 通信機構の設計(2012), SACSIS2012(2) 今出広明 et al., 大規模計算環境のためのランク配置最適化手法 RMATT(2011),
SACSIS2011(3) T. Adachi et al., The design of ultra scalable MPI collective communication on
the K computer(2012), ISC'12

Tofu
Tofu

MPI

MPI
2012 10 24

Copyright 2012 FUJITSU LIMITED

1

Copyright 2012 FUJITSU LIMITED: SPARC64TM VIIIfx(45nm)
1

8
1

100
24
IO

:ICC
6

(Tofu)

10

:

Tofu

(10PFlops: 864

4

)

ICC
LSI

(

250
(8

2

Copyright 2012 FUJITSU LIMITED

Tofu

)
)

3

Copyright 2012 FUJITSU LIMITEDInterConnect Controller ICC Chip

SPARC64TM VIIIfx
Torus fu

RDMA

-Torus 3D-Torus = 6D-Torus
6

PCI Express
312.5 MHzX, Y, Z, A, B, C
32, 32, 32, 2, 3, 2

65nm ASIC: X, Z, B:Z=1

SPARC64TM VIIIfx
CPUC: Y, A, C
16

IO

18.2mm x 18.1mm
48M
SRAM 12Mbit
128

:Z=0

Y
Y
X
B
X

InterConnect
Controller
(ICC)

5GB/s
140GB/sXYZ

A Z Z B
4

4
10

Virtual Cut-ThroughABC
Copyright 2012 FUJITSU LIMITED

5~100nsCopyright 2012 FUJITSU LIMITED

ICC

Tofu
10 XYZ
RDMA

5GB/s

6

ABC 4) 10
4

RDMA

0

RDMA

1

RDMA

2

RDMA

3

6
4
link 0link 1link 2link 3link 4link 5link 6link 7link 8link 9

CPU

X

4
XYZ
ABC
LSI

Y
Y

A

B
ReduceAllreduce
64

CPU: AND, OR, XOR, MAX, SUM
160

10
C

: SUM

ICC

7.8
6

Z

Y

Copyright 2012 FUJITSU LIMITED

7

Tofu

Copyright 2012 FUJITSU LIMITED

6
CA

CPU

Node: 1 CPU + 1 ICC (Tofu Chip)
Topology: 6D Mesh/Torus
Mesh: Y, A, C / Torus: X, Z, B
SW

A{2} x C{2} on a board
B{3} x Z{17} in two racks

B

SW
SW

48 boards: Z=1-16, compute nodes
12 I/O boards: Z=0, I/O nodes

Z

2

SW

Tofu

2

(6

㻺㼡㼙㼎㼑
r 㼛㼒㻯㼍
㼎
㼘
㼑㼟
㻦
㻞㻜㻜㻷

Z

SW
SW

3

/

Fat-Tree

)

3

Tofu

Fat-tree

X
Y

**
*

㻯㼛㼡r㼠
㼑㼟
㼥㼛㼒㻾㻵
㻷㻱㻺

(

*

8

**TofuCopyright 2012 FUJITSU LIMITED

6

)

/

9

Tofu

Copyright 2012 FUJITSU LIMITED

2,3

vs.Tofu

Tofu XYZ ABC
1

2

3
2

3

X=x+a, Y=y+b, Z=z+c
X=x+a+b, Y=y+z+c 1
X=x+y+z+a+b+c

:
Tofu:
2,3

Tofu
a

X

(

)

X

(

)

x

(

)

a=1
2

1

1

2

a=0

:

㼄

Tofu:

㼆

2,3

㼅 㼅㻩㻔
㼥
㻗㼎
㻕

3

2
㻮1
0

Tofu
b

Y

(

0

1

2

3

4

5

16

15

3

4

5

8

b=1

17

14

13

12

6

7

b=0

0

1

2

11

10

9

10

)

Y

(

)

y

(

)

b=2

Copyright 2012 FUJITSU LIMITED

11

Copyright 2012 FUJITSU LIMITED

12
x y z

(2x3x2)a c b(b) (c) (a) x yz a c bMPI

12Copyright 2012 FUJITSU LIMITED

13

Copyright 2012 FUJITSU LIMITED

MPI

MPI
MPI (Message Passing Interface)
HPC

8
1
1
(OpenMP, VISIMPACT)

MPI-2.1
MPI
1

1

MPI

MPI-IO1
MPI-3.0

2012

9 21

Non-blocking Collective
Neighborhood Collective: Blocking Non-blocking,One-sided: RDMA
Get,Put,Atomic
14

Copyright 2012 FUJITSU LIMITED

15

MPI

Copyright 2012 FUJITSU LIMITED

MPI

8
MPI

8

8

11.6GB

SimFold

8ModylasAllgatherv, Allreduce

GAMESS

Allreduce, Bcast

RSDFT

Allreduce, Bcast,

LANS

3

NICAM

Send-Recv(

LatticeQCD

6HPL

Allgather, Bcast, Send-Recv

NPB-FT

Alltoall

Tofu,Reduce,Reduce_Scatter, Allgather

Bcast, Allreduce, Alltoall, Allgather

19)

Send-Recv1

1

10% 1.6GB
16

Copyright 2012 FUJITSU LIMITED

17

Copyright 2012 FUJITSU LIMITED

MPI
Eager

Rendezvous
RDMA WRITE(Eager18)

Copyright 2012 FUJITSU LIMITED

19

Copyright 2012 FUJITSU LIMITED

MPI
8

11.6GB(RSDFT, LANS, LQCD)

Open MPI
1<>

processD
processA

NICAM

145KB
145B
256
4

processB
processC

19

(

)

㻝㻝
㻚
㻢㻳㻮

8

20

Copyright 2012 FUJITSU LIMITED21

Copyright 2012 FUJITSU LIMITED

Alltoall
1.0E+11

AlltoAll(MAX=1024)

㻥㻞㻝
㻢
㻝㻟
㻳㻮

AlltoAll(Only Fast Model)

1M(16Kx64
2K(128 x16

)

㻝㻟
㻘
㻣㻜
㻜

AlltoAll(Only Resouce Saving Model)
Init/Finalize

1.0E+10)

Rendezvous

㻥㻞㻝
㻢
㻝㻚
㻡㻳㻮

1.0E+09~128~13.7K

13.7K

1M(16K

x64

)

Eager

Eager

Rendezvous

2K(128

x16

)

Eager

Rendezvous

Rendezvous

15
1024

~

㻥㻞㻝
㻢
㻜㻚
㻝
㻠㻳㻮

1.0E+08

㻵
㼚㼕
㼠
㻛㻲㼕
㼚㼍
㼘㼕
㼦
㼑
㻥㻞㻝
㻢
㻜㻚
㻜㻣㻳㻮

1024
(1G

1.0E+07

)

0

2000

4000

6000

8000

10000

# of Nodes
22

Copyright 2012 FUJITSU LIMITED

23

Copyright 2012 FUJITSU LIMITED

10PF
10PF

10%

1.0

Init-Finalize
3D
Alltoall
Alltoall 1024

0.0
010,00020,00030,00040,00050,00060,000

Alltoall: Simple
Spread70,00080,00090,000# of nodes
24

Tofu

(4TNI,

Copyright 2012 FUJITSU LIMITED25

MPI_Bcast(

)

)

1

TNI1

TNI0

Copyright 2012 FUJITSU LIMITED

TNI

26

2

Copyright 2012 FUJITSU LIMITED

2

27

3

Copyright 2012 FUJITSU LIMITED(Trinaryx3)24

22in2out

2in2out3(Trinaryx3)
Tofu

3

TNI

1/3

3

Step2

3out

Step3

Step4

Step1

28

Copyright 2012 FUJITSU LIMITED

29

Copyright 2012 FUJITSU LIMITEDBcast(MPI_Alltoall)

Trinaryx3

11

MPI_Alltoall

(

)

60%

( Tofu30

Copyright 2012 FUJITSU LIMITED

)

31

1

(8x6x32 alltoall(

Copyright 2012 FUJITSU LIMITED33(MPI_Allreduce)(

Copyright 2012 FUJITSU LIMITED(MPI_Allgather))

Trinaryx3 Reduce Trinaryx3 Bcast
Reduce)-1)

32

MPI_Allreduce

Copyright 2012 FUJITSU LIMITEDMPI_Allgather
reduce

4

Bcast

(

)

4

1D ring

1D ring
bcast

Reducephase1 1
phase2phase3

+
Bcast

34

Copyright 2012 FUJITSU LIMITED

35

Copyright 2012 FUJITSU LIMITED(MPI_Barrier)

10PF

Tofu

(TBI) MPI

Tofu

(MPI_Bcast)

10PF

(

114

132

)

+TBI

Allreduce(MPI_SUM

Bcast (6) Bcast

9

)

11

7.2
Barrier / Allreduce
100

Barrier (TBI)
Barrier (soft)

90

Allreduce (TBI)
Allreduce (soft)

80

1048x54x32

9

48x48x328

48x42x32
48x36x32

787.9us

70

6

60

5
9,216

50

4

40

3

30

2

20

1

10

12.2us

0

0
0

1,000

2,000

3,000

4,000

5,000

6,000

7,000

8,000

1.E+0

9,000 10,000

1.E+1

1.E+2

1.E+3

1.E+4

1.E+5

1.E+6

1.E+7

1.E+8

1.E+9

(B)

nodes (48x6xZ)
36

10PF

Copyright 2012 FUJITSU LIMITED(MPI_Allreduce)

10PF(
Allreduce (6

114
9

37

10PF
132

)

Copyright 2012 FUJITSU LIMITED(MPI_Allgather)

10PF
Allgather (6

)

9

)

13

7
48x54x32

6

48x48x32

48x54x32

11

48x48x321048x42x32

512

48x42x32

948x36x32

48x36x32

8

4

7
6

3

5
4

2

3
2

1

1
0

0
1.E+0

1.E+1

1.E+2

1.E+3

1.E+4

1.E+5

1.E+6

1.E+7

1.E+8

1.E+9

1.E+0

1.E+1

1.E+2

1.E+3

38

1.E+4

Copyright 2012 FUJITSU LIMITED

39

Linpack()

RSDFT(

Copyright 2012 FUJITSU LIMITED

Bcast
Allgatherv

10.51Pflops
93.17%

Allreduce(

3.08Pflops
43.63%)(HPCC

34.7Tflops*))
HPCC(Global FFT)()
㻿㻯㻝㻝㼀㻻㻼㻡㻜㻜
㻏㻝

Alltoall㻿㻯㻝㻝㻳㼛rd㼛㼚㻮㼑㼘
㼘
㻼r㼕
㼦
㼑

* : using FFTE
40

Copyright 2012 FUJITSU LIMITED

1.E+5

(B)

(B)

41

㻴㻼㻯㻯㼔㼍
㼘㼘
㼑㼚
㼓
㼑㻭㼣㼍
rd㼟

developed by Prof . Takahashi of Tsukuba Univ.
Copyright 2012 FUJITSU LIMITEDRSDFT

Allreduce
Allreduce

2011 Gordon Bell Award*

O riginal

Modified for K com puter

19%
32%

Com putation

Com m unication

Com putation(8,000 atom s,256 nodes)

allreduce

Com m unication(107,292 atom s,55,296 nodes)::3.08Pflops (43.6% efficiency )3.2G B/s.

*:
Proceedings of 2011 InternationalConference for H igh Perform ance Com puting,N etw orking,Storage and Analy sis,
Article N o. 1 ,First-principles calculations of electron states of a silicon nanow ire w ith 100,000 atom s on the K com puter

Courtesy of RIKEN

42

Copyright 2012 FUJITSU LIMITED

0

1

2

8

6

2

3

4

5

1

3

0

6

7

8

5

7

4

43

Copyright 2012 FUJITSU LIMITED

45

Copyright 2012 FUJITSU LIMITED

2

N

N!

SA: Simulated Annealing
44

Copyright 2012 FUJITSU LIMITED

RMATT

䢿

SA

NAS Parallel Benchmark

RMATT(Rank Mapping Automatic Tuning Tool)

Case 1䣔M䣃䣖䣖

NPROCS=1024, CLASS=B
1䢰
33

2D Torus 32x32

1䢰
24

䢹䢧

䣔M䣃䣖䣖
32x32x16

MPI

RMATT

0 (0,0,0)
1 (0,0,1)
2 (0,1,0)(machinefile

)

MPI
RMATT

x,y
MPI

Case 2䣔M䣃䣖䣖

NPROCS=8192, CLASS=D

10䢰䢻42D Torus 128x64
46

Copyright 2012 FUJITSU LIMITED

47

䢻䢰
䢻䢺

䢸4䢧
䢻䢧
Copyright 2012 FUJITSU LIMITED

MPI

8
1

48

Copyright 2012 FUJITSU LIMITED

49

Copyright 2012 FUJITSU LIMITED

科 学 技 術 計 算 分 科 会

選 出

HPC フォーラム 2012 より

Algorithms and Softwarein the Post-Petascale Era

William D. Gropp(University of Illinois)

Extrapolation is Risky
Algorithms and Software in
the Post-Petascale Era• 1989 – T – 23 years
 Intel introduces 486DX
 Eugene Brooks writes “Attack of the Killer

William Gropp
www.cs.illinois.edu/~wgropp

Micros”
 4 years before TOP500
 Top systems at about 2 GF Peak• 1999 – T – 13 years
 NVIDIA introduces its GPU (GeForce 256)
• Programming GPUs still a challenge 13 years later Top system – ASCI Red, 9632 cores, 3.2 TF

Peak (about 3 GPUs in 2012)

 MPI is 7 years old

HPC Today

2

HPC in 2011
• Sustained PF systems• High(est)-End systems
 1 PF (1015 Ops/s) achieved on a few “peak friendly”

applications
 Much worry about scalability, how we’re going to get to an
ExaFLOPS Systems are all oversubscribed• DOE INCITE awarded almost 900M processor hours in 2009;
1600M-1700M hours in 2010-2012; (big jump planned in 2013
– over 5B hours)• NSF PRAC awards for Blue Waters similarly competitive• Widespread use of clusters, many with accelerators;
cloud computing services
 These are transforming the low and midrange• Laptops (far) more powerful than the supercomputers I
used as a graduate student Kei Computer (Fujitsu) at RIKEN, Kobe, Japan (2011)
 “Sequoia” Blue Gene/Q at LLNL NSF Track 1 “Blue Waters” at Illinois
 Undoubtedly others (China, … )• Still programmed with MPI and MPI+other (e.g.,
MPI+OpenMP or MPI+OpenCL/CUDA or MPI+OpenACC)
 But in many cases using toolkits, libraries, and otherapproaches• And not so bad – applications will be able to run when thesystem is turned on Replacing MPI will require some compromise – e.g.,domain specific (higher-level but less general)• Lots of evidence that fully automatic solutions won’t work

3

4

End of an Era

The Post-Moore Era• IN THE LONG TERM (~2017 THROUGH 2024)“While power consumption is an urgent
challenge, its leakage or static component will
become a major industry crisis in the longterm, threatening the survival of CMOS
technology itself, just as bipolar technology
was threatened and eventually disposed of
decades ago.” [ITRS 2009]• Unlike the situation at the end of the bipolarera, no technology is waiting in the wings.

5• Scaling is ending Voltage scaling ended in 2004 (leakage current)
 Feature scaling will end in 202x (not enough atoms)
 Scaling rate will slow down in the next few years
 Continued scaling in the next decade will need asequence of (small) miracles (new materials, newstructures, new manufacturing technologies)☛Compute Efficiency becomes a paramountconcern
 More computations per joule More computations per transistor

6

HPC in 2020-2023
2018-2020• Exascale systems are likely to have
 Extreme power constraints, leading to• Clock Rates similar to today’s systems• A wide-diversity of simple computing elements (simple for
hardware but complex for software)• Memory per core and per FLOP will be much smaller• Moving data anywhere will be expensive (time and power)
 Faults that will need to be detected and managed• Some detection may be the job of the programmer, as
hardware detection takes power Extreme scalability and performance irregularity• Performance will require enormous concurrency• Performance is likely to be variable Simple, static decompositions will not scale A need for latency tolerant algorithms andprogramming

Algorithms and Applications
Will Change• Applications need to become more
dynamic, more integrated• System software must work withapplication:
 Code complexity (Autotuning) Dynamic resources (no simple PGAS) Latency hiding (Nonblocking algorithms,

interfaces (including futures)) Resource sharing (more performance

information, performance asserts, runtime
coordination)• Memory, processors will be 100s to 10000s of cycles
away. Waiting for operations to complete will cripple
performance
7

8

How Do We Make Effective
Use of These Systems?

Make the Compiler Better• Better use of our existing systems
 Blue Waters will provide a sustained PF, but thattypically requires ~10PF peak

• Improve node performance
 Make the compiler better Give better code to the compiler Get realistic with algorithms/data structures• Improve parallel performance/scalability• Improve productivity of applications
 Better tools and interoperable languages, not a (single)• It remains the case that mostcompilers cannot compete withhand-tuned or autotuned code on
simple code Just look at dense matrix-matrix

multiplication or matrix transpose Try it yourself!

• Matrix multiply on my laptop:• N=100 (in cache): 1818 MF (1.1ms)• N=1000 (not): 335 MF (6s)new programming language• Improve algorithms
 Optimize for the real issues – data movement, power,

resilience, …
9

10

How Good are Compilers at
Vectorizing Codes?

Media Bench II Applications
Appl

XLC

ICC

GCC

XLC

Automatic

S. Maleki, Y. Gao, T. Wong, M. Garzarán, and D. Padua. An Evaluation of Vectorizing
Compilers. PACT 2011.
11

ICC

GCC

ManualJPEG Enc

-

1.33

-

1.39

2.13

1.57

JEPG Dec

-

-

-

-

1.14

1.13

H263 Enc

-

-

-

1.25

2.28

2.06

H263 Dec

-

-

-

1.31

1.45

-

MPEG2 Enc

-

-

-

1.06

1.96

2.43

MPEG2 Dec

-

-

1.15

1.37

1.45

1.55

MPEG4 Enc

-

-

-

1.44

1.81

1.74

MPEG4 Dec

-

-

-

1.12

-

1.18

Table shows whole program speedups measured
against unvectorized application
S. Maleki, Y. Gao, T. Wong, M. Garzarán, and D. Padua. An Evaluation of Vectorizing Compilers. PACT 2011.
12

How Do We Change This?
•

Test compiler against “equivalent” code (e.g., best hand-tuned or
autotuned code that performs the same computation, under someinterpretation or “same”)•

As part of the Blue Waters project, Padua, Garzaran, Maleki aredeveloping a test suite that evaluates how the compiler does withsuch equivalent codeIn a perfect world, the compiler would provide the same, excellent
performance for all equivalent versions

Working with vendors to improve the compiler
Identify necessary transformations
Identify opportunities for better interaction with the programmer tofacilitate manual intervention.
 Main focus has been on code generation for vector extensions
 Result is a compiler whose realized performance is less sensitive to different
expression of code and therefore closer to that of the best hand-tunedcode.
 Just by improving automatic vectorization, loop speedups of more than 5
have been observed on the Power 7.




•

But this is a long-term projectGive “Better” Code to theCompiler• Augmenting current programming
models and languages to exploit
advanced techniques for
performance optimization (i.e.,
autotuning)• Not a new idea, and some tools
already do this.
• But how can these approaches
become part of the mainstream
development?

What can we do in the meantime?
13

14

How Can Autotuning Tools Fit
Into Application Development?

Application Needs Include• In the short run, just need effective
mechanisms to replace user code with
tuned code Manual extraction of code, specification of

specific collections of code transformations• But this produces at least two versionsof the code (tuned (for a particulararchitecture and problem choice) anduntuned). And there are other issues.
• What does an application want (what isthe Dream)?

• Code must be portable• Code must be persistent• Code must permit (and encourage)
experimentation• Code must be maintainable• Code must be correct• Code must be faster

15

16

Implications of These
Requirements

Application-Relevant
Abstractions• Portable - augment existing language. Either use
pragmas/comments or extremely portable precompiler Best if the tool that performs all of these steps looks like just likethe compiler, for integration with build process• Persistent Keep original and transformed code around: Golden Copy

• Maintainable Let user work with original code and ensure changes automatically

update tuned code

• Correct
 Do whatever the application developer needs to believe that the

tuned code is correct• In the end, this will require running some comparison tests• Faster Must be able to interchange tuning tools - pick the best tool foreach part of the code

• Language for interfacing with autotuning must convey
concepts that are meaningful to the application
programmer• Wrong: unroll by 5
 Though could be ok for performance expert, and somecompilers already provide pragmas for specifictransformations• Right (maybe): Performance precious, typical loop countbetween 100 and 10000, even, not power of 2
• Middle ground: Apply unroll, align, SIMD
transformations and tune
• We need work at developing higher-level, performanceoriented languages or language extensions
 This would be the “good” future
 Early steps include TCE, Orio, Spiral, …

 No captive interfaces
 Extensibility - a clean way to add new tools, transformations,

properties, …

17

18

Better Algorithms and Data
Structures• Autotuning only offers the best
performance with the given data
structure and algorithm

Sparse Matrix-Vector Multiply
Barriers to faster code

• Processors include hardware to address
performance challenges
 “Vector” function units Memory latency hiding/prefetch
 Atomic update features for shared memory
 Etc.
19

What Does This Mean For
You?

2.00

Performance Ratio That’s a big constraint• “Standard” formats
such as CSR do notmeet requirementsfor prefetch or
vectorization• Modest changes todata structure
enable both
vectorization,
prefetch, for 2080% improvement
on P7

1.80
1.60
1.40

SCSR2
SCSR4
VSCSR
-2
VSCSR
-4

1.20
1.00
0.80

Prefetch results in Optimizing Sparse Data
Structures for Matrix Vector Multiply
http://hpc.sagepub.com/content/25/1/115
20Is It Communication Avoiding Or
Minimum Solution Time?

• It is time to rethink data structures andalgorithms to match the realities of memoryarchitecture We have results for x86 where the benefit is smaller

but still significant

 Better match of algorithms to prefetch hardware isnecessary to overcome memory performance
barriers• Similar issues come up with heterogeneous
processing elements (someone needs todesign for memory motion and concurrent andnonblocking data motion)• Example: non minimum collective
algorithms• Work of Paul Sack; see “Faster
topology-aware collective
algorithms through non-minimalcommunication”, PPoPP 2012
• Lesson: minimum communicationneed not be optimal

21

22

Allgather

Problem: Recursive-doubling• No congestion model:

1

2

3

4

 T=(lgP)α + n(P-1)β• Congestion on torus:

Input T≈(lgP)α + (5/24)nP4/3β

Output• Congestion on Clos network: T≈(lgP)α + (nP/μ)β• Solution approach: move smallest
amounts of data the longest distance23

24

Allgather: recursive halving

Allgather: recursive halving

a

b

c

d

ac

bd

ac

bd

e

f

g

h

eg

fh

eg

fh

i

j

k

l

ik

jl

ik

jl

m

n

o

p

mo

np

mo

np

25

26

Allgather: recursive halving

Allgather: recursive halving

acik

bdjl

acik

bdjl

acikbdjl

acikbdjl

acikbdjl

acikbdjl

egmo

fhnp

egmo

fhnp

egmofhnp

egmofhnp

egmofhnp

egmofhnp

acik

bdjl

acik

bdjl

acikbdjl

acikbdjl

acikbdjl

acikbdjl

egmo

fhnp

egmo

fhnp

egmofhnp

egmofhnp

egmofhnp

egmofhnp27

28

Allgather: recursive halving

New Problem: Data
Misordered• Solution: shuffle input data

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

acikbdjl
egmofhnp

 Could shuffle at end (redundant

work; all processes shuffle) Could use non-contiguous datamoves
 Shuffle data on network…

T=(lg P)α + (7/6)nPβ2930

Solution: Input shuffle

Solution: Input shuffle

a

b

c

d

a

e

b

f

e

f

g

h

i

m

j

n

i

j

k

l

c

g

d

h

m

n

o

p

k

o

l

p

31

32

Solution: Input shuffle

Solution: Input shuffle

a

e

b

f

ab

ef

ab

ef

i

m

j

n

ij

mn

ij

mn

c

g

d

h

cd

gh

cd

gh

k

o

l

p

kl

op

kl

op

33

34

Solution: Input shuffle

Evaluation:
Intrepid BlueGene/P at ANL• 40k-node systemabcdefghabcdefghabcdefghabcdefgh

ijklmnop

ijklmnop

ijklmnop

ijklmnop

abcdefghabcdefghabcdefghabcdefgh

ijklmnop

ijklmnop

ijklmnop

ijklmnop Each is 4 x 850 MHz PowerPC 450• 512+ nodes is 3d torus; fewer is
3d mesh• xlc -O4
• 375 MB/s delivered per link
 7% penalty using all 6 links both

T=(1+lgP) α + (7/6)nPβ
T≈(lgP)α35 + (7/6)nPβ

ways

36

Allgather performance

37

Performance on a Node
• Nodes are SMPs
 You have this problem on anything(even laptops)• Tuning issues include the usual Getting good performance out of thecompiler (often means adapting to
the memory hierarchy)• New (SMP) issues include Sharing the SMP with other processes
 Sharing the memory system

Notes on Allgather• Bucket algorithm (not describedhere) exploits multiple
communication engines on BG• Analysis shows performance near
optimal• Alternative to reorder data step isin memory move; analysis shows
similar performance andmeasurements show reorder step
faster on tested systems
38

New (?) Wrinkle – Avoiding
Jitter• Jitter here means the variation in timemeasured when running identicalcomputations
 Caused by other computations, e.g., an OSinterrupt to handle a network event or
runtime library servicing a communication
or I/O request• This problem is in some ways lessserious on HPC platform, as the OS andruntime services are tuned to minimize
impact However, cannot be eliminated entirely

39

40

Sharing an SMP

Happy Medium Scheduling• Having many cores available
makes everyone think thatthey can use them to solve
other problems (“no one
would use all of them all of
the time”)• However, compute-boundscientific calculations areoften written as if all computeresources are owned by theapplication• Such static scheduling leads
to performance loss
• Pure dynamic scheduling adds
overhead, but is better• Careful mixed strategies areeven better• Thanks to Vivek Kale

Performance irregularities introduce loadimbalance.
Pure dynamic has significant overhead; purestatic too much imbalance.
Solution: combined static and dynamicscheduling
Communication Avoiding LU factorization(CALU) algorithm, S. Donfack, L .Grigori, V.
Kale, WG, IPDPS ‘12

Scary Consequence: Static
data decompositions will notwork at scale.
Corollary: programming
models with static task
models will not work at scale
41

42

Synchronization and OS Noise

A Noisy Example –Dissemination Barrier• “Characterizing the Influence ofSystem Noise on Large-Scale
Applications by Simulation,”
Torsten Hoefler, Timo Schneider,
Andrew Lumsdaine
 Best Paper, SC10

• Next 3 slides based on this talk…

• Process 4 is delayed
 Noise propagates “wildly” (of course

deterministic)

43

44

Single Collective Operations
and Noise

The problem isblocking operations

Legend:deterministic

3rdquartilemedian
2ndquartileoutliers

outliers• Simple, data-parallel algorithms easy toreason about but inefficient True for decades, but ignored (memory)• One solution: fully asynchronous
methods
 Very attractive, yet efficiency is low andthere are good reasons for that

 Blocking can be due to fully collective (e.g.,• 1 Byte, Dissemination, regular noise,1000 Hz, 100 μs

Allreduce) or neighbor communications(halo exchange) Can we save methods that involve global,
synchronizing operations?

45

46

Saving Allreduce

The Conjugate Gradient
Algorithm• One common suggestion is to avoid using
Allreduce
 But algorithms with dot products are among the bestknown Can sometimes aggregate the data to reduce thenumber of separate Allreduce operations
 But better is to reduce the impact of thesynchronization by hiding the Allreduce behind otheroperations (in MPI, using MPI_Iallreduce)• We can adapt CG to nonblocking Allreducewith some added floating point (but perhaps
little time cost)47• While (not converged)
niters += 1;s= A * p;t
= p' *s;alpha = gmma / t;x
= x + alpha * p;r
= r - alpha * s;
if rnorm2 < tol2 ; break ; endz= M * r;gmmaNew = r' * z;beta = gmmaNew / gmma;
gmma = gmmaNew;p= z + beta * p;
end
48The Conjugate Gradient
Algorithm• While (not converged)
niters += 1;s= A * p;t
= p' *s;alpha = gmma / t;x
= x + alpha * p;r
= r - alpha * s;
if rnorm2 < tol2 ; break ; endz= M * r;gmmaNew = r' * z;beta = gmmaNew / gmma;
gmma = gmmaNew;p= z + beta * p;end

CG Reconsidered• By reordering operations, nonblocking
dot products (MPI_Iallreduce in MPI-3)
can be overlapped with other operations• Trades extra local work for overlapped
communication
 On a pure floating point basis, thenonblocking version requires 2 more
DAXPY operations
 A closer analysis shows that some
operations can be merged• More work does not imply more time

49

50

What’s Different at Peta/Exascale

Preparing for the Next
Generation of HPC Systems• Performance Focus

 Only a little – basically, the resource is expensive, so a

premium placed on making good use of resource Quite a bit – node is more complex, has more featuresthat must be exploited• Scalability

 Solutions that work at 100-1000 way often inefficient at

100,000-way Some algorithms scale well

• Explicit time marching in 3D
 Some don’t• Direct implicit methods
 Some scale well for a while• FFTs (communication volume in Alltoall) Load balance, latency are critical issues• Fault Tolerance becoming important Now: Reduce time spent in checkpoints
 Soon: Lightweight recovery from transient errors
51• Better use of existing resources Performance-oriented programming Dynamic management of resources at all levels
 Embrace hybrid programming models (you havealready if you use SSE/VSX/OpenMP/…)• Focus on results Adapt to available network bandwidth and latency
 Exploit I/O capability (available space grew fasterthan processor performance!)• Prepare for the future Fault tolerance
 Hybrid processor architectures
 Latency tolerant algorithms
 Data-driven systems

Recommended Reading• Bit reversal on uniprocessors (Alan Karp, SIAM
Review, 1996)• Achieving high sustained performance in an
unstructured mesh CFD application (W. K.
Anderson, W. D. Gropp, D. K. Kaushik, D. E.
Keyes, B. F. Smith, Proceedings ofSupercomputing, 1999)• Experimental Analysis of Algorithms(Catherine McGeoch, Notices of the American
Mathematical Society, March 2001)• Reflections on the Memory Wall (Sally McKee,
ACM Conference on Computing Frontiers,
2004)
53

52

Thanks• Torsten Hoefler• Abhinav Bhatele

 Performance modeling lead,Blue Waters; MPI datatype

Process/node mapping• Elena Caraba• David Padua, Maria
Garzaran, Saeed Maleki Nonblocking Allreduce in CG• Van Bui

 Compiler vectorization

 Performance model-based• Dahai Guo

evaluation of programming
models Streamed format exploiting• Funding provided by:prefetch, vectorization, GPU• Vivek Kale

 Blue Waters project (State of

Illinois and the University of
Illinois)

 SMP work partitioning• Hormozd Gahvari Department of Energy, Office AMG application modelingof Science• Marc Snir and William
Kramer National Science Foundation

 Performance model

advocates
54