大規模スパコンにおけるHPCとビッグデータ・AIの統合

理化学研究所 計算科学研究センターセンター長 松岡 聡
トーゴシンポジウム 講演
20181005

1

TSUBAME3.0

東工大スパコンTSUBAMEシリーズの歴史
2010年
最先端スパコン・15年間で約一万倍の性能向上
京コンピュータ
2000年 地球シミュレータ
エコシステムの最重視
11.4ペタフロップス
松岡GSIC 40テラフロップス
2002年

に赴任

1988-2006
専用スパコン

2006年 TSUBAME1.0
80テラフロップス
アジア一位・世界7位

2008年 TSUBAME1.2
170テラフロップス
世界初大規模GPUスパコン

2010年 TSUBAME2.0
2.4ペタフロップス・世界4位
運用スパコングリーン世界一
ゴードンベル賞

2002年 TSUBAME0
1.3 テラフロップス
日本スパコンセンター初の大規模クラスタ計算機 汎用CPUやメニーコアプロセッサ(GPU)活用・光ネットワーク・

不揮発性メモリ・高効率電力制御や水冷など

2013年TSUBAME2.5
4108機のGPUアップグレード
5.7ペタフロップス・日本2位
単精度演算日本一 17.1ペタ
1000万CUDA コア級

2013年TSUBAME-KFC
TSUBAME3技術実験機
液浸・高温冷却技術
スパコングリーン世界一
2015年AIスパコンプロト(KFC/DL)

2017年 TSUBAME3.0
12.1ペタフロップス (AI性能47.2ペタ)スパコングリーン世界一(大規模実用スパコンでは初)ビッグデータ・ＡＩとHPCの統合
1000万CUDA コア級

• 遺伝研・DDBJスパコ
ンはTSUBAME2をベースに遺伝研チームが設計– 松岡もヘルプ• ゲノム処理を加速す
るように種々のパラ
メタを特化– メモリ・I/O強化– Assy用にFatノード– GPUは少なめ

2018 4/1 より理研計算科学センタ―長に計算の 計算による 計算のための科学
理化学研究所 計算科学研究センター（R-CCS）

計算の科学

計算による科学

生命科学、工学、気象・気候、防「京」やポスト「京」を支えるプロ
災・減災など私たちの生活に直結
グラミング手法・ソフトウェア・運し、国民の関心事の高い最先端の用技術や、ビッグデータ・AIなどへ研究開発に欠くことのできない、の対応を実現する手法など、様々な基礎研究や応用研究を「京」やポ技術のベースとなる高性能計算の本
スト「京」を用いて推進質に関する研究を推進
シナジー、融合

国内外の各機関、大学、企業の研究者との協力を更に拡大

世界中からトップ研究センターと認知される
国際的な「計算のための科学」の知の拠点

4

R-CCSの特筆すべき業績等

災害被害シミュレーション

交通ネットワーク（社会科学）復旧シミュレーション

単純計算の速度（TOP500） 9倍
1位
10位
京

産業利用ソフトの速度

ビッグデータ処理速度
2/3
1位

2位

SunwayTaihuLight（中国）

HPCG 性能

都市圏を丸ごとモデル化

ライフライン「京」の性能を引き出す独自のアルゴリズムを研究開発した成果として、大規模グラフ解析に関するスパコン性能指標Graph500で、平成27年6月から平成29年11月まで6期連続（通算7期）で世界1位を獲得。産業利用
ソフトで用いる計算手法の処理速度に関するスパコン性能指標HPCGで、
平成28年11月から平成29年11月まで3期連続で世界1位を獲得。
Graph500 性能「京」全体で開発ソフトウェアGAMERAを実行し、従来の1,000倍以上の規模の地盤振動問題を解くことで、高性能計算技術の世界最高峰の国際会議
SC14,SC15でゴードン・ベル賞ファイナリストに選出され、SC16,SC17で最優
秀ポスター賞を受賞。

Graph500,HPCGで世界1位を獲得

LINPACK 性能

次世代地震被害予測システムのコア技術で受賞

4/5
1位

5位

※平成29年度11月時点

スパコン向けソフト開発を大幅に容易に「3D降水ナウキャスト手法」を用いた降水予報

並列化や最適化のプログラミングを自動化することで大規模スパコン向けシミュレーションソフトの開発コストを大幅削減。
高性能計算技術の世界最高峰の国際会議SC16で最優秀論文賞を受賞。

解像度100mで30秒ごとに新しい観測データを取り込んでリアルタイムに予測を実行するシステム（3D降水ナウキャスト手法）を構築し、世界初となる30秒
更新10分後までの降水予報を開始した。

ソフト開発コストの大幅削減等が期待される。

将来、超高速かつ超高精細な天気予報が可能になると期待される。

5

分子動力学プログラムGENESIS• 2010年から主にスパコン「京」の効率的な利用のために、粒子系生物物理研究チームはGENESIS という新しい分子動力学プログラムの開発を始めた。
• Jaewoon Jung, Takaharu Mori, et al. WIREs CMS 5, 310-323 (2015).
• Chigusa Kobayashi, Jaewoon Jung et al. J. Comp. Chem. (2017)• GENESISの高速化と並列化• Inverse Lookup Table: Jung et al. JCC 34, 2414-2420 (2013)• Midpoint Cell Method: Jung et al. JCC 35, 1064-1072 (2014)• Volumetric 3D FFT: Jung et al. CPC 200, 57-65 (2016)• GPU parallelization: Jung et al. JCTC 12, 4947-4958 (2016)• ポスト「京」重点アプリの一つとして、コデザインを実施中

New Version (2.0.0) 未公開

GENESISのベンチマーク
Force Field: AMBER FF
MD condition: Cutoff= 8Å, NVE, Single precision
Computer: RIKEN BigWaterFall(Intel Xeon Gold 6148 (2.4GHz): 20 core, 2CPU/node1680CPU) )DHFR (23,558 atoms)

Cellulose (408,609 atoms)
128

128

64

64
32

128

Hydrogen mass
repartitioning∆t=4fs64

ns/day32

256

ns/dayns/day512

STMV (1,067,095 atoms)

16
8

RESPA (long = 2)∆t=2fs

4
2

1

2

4

8

# of nodes (x 40 cores)

16
8
Hydrogen mass
repartitioning∆t=4fs

4
2
1

1

2

4

8

16 32 64

# of nodes (x40 cores)

1

2

4

8

16 32 64

# of nodes (x 40 cores)バクテリア細胞質の全原子MD計算
1億以上の原子を含む世界で初めてのバクテリア細胞質の全原子MD計
算を実現
Yu, Mori, Harada, Ando, Jung, Sugita, Feig, eLife (2016) 5, e19274

In collaboration with M. Feig (MSU)

バクテリア細胞質の全原子MD計算
RIKEN Channel (Youtube) で公開中

JPN：https://youtu.be/2JU2LjPDrQY

ENG：https://youtu.be/5JcFgj2gHx8

9分子動力学計算の創薬応用（従来の方法）蛋白質と薬剤候補化合物の剛体ドッキング（これから）蛋白質と薬剤候補化合物を含む分子動力学

Binding site

”Key”
Ligand“Lock”
Protein receptor

立体構造の相補性

立体構造変化、分子間相互作用

今後：スパコンとビッグデータ基盤の統合
JST-CREST EBD-Extreme Big Data (2013-2018)

トップクラスのス
パコンによって、現状のクラウド比
でどの程度将来のビッグデータア
プリを加速できるのか？

大規模グラフ処理
社会シミュレーション

大規模メタ
ジェノミックス

Co-Design
EBD Bag

大規模センサー系
計算とデータ処理のリアルタイム統合

Co-Design Co-Design

EBDシステムソフトウェア
ソフトウェアスタック

Cartesian Plane
KV
S

KV
S

Graph Store

NVM/Fla
NVM/Flas
2Tbps HBM
NVM/Flashh
4~6HBM Channels NVM/Flas
NVM/Fla
NVM/Flasshh
1.5TB/s DRAM & h
sh
DRAM
DRAM
NVM BW
DRAM
DRAM
DRAM
DRAM
Low
Low
High Powered
30PB/sI/O BW Possible
Main CPU
Power
Power
1
Yottabyte
/
Year
CPU
CPU
TSV Interposer

KVS

EBD KVS

次世代の統合インフラ
PCB

統合アーキテクチャ
大容量不揮発性メモリ、超広帯域光ネットワーク等

クラウドIDC
低バンド幅・低効率

アーキテク
チャ、アルゴリ
ズム、システム
ソフトウェア、性能モデリングなど
新たなアーキやデバイスの活用は?
スパコン
計算バッチ処理中心・
低柔軟性・高コスト

No. of edges

K computer: 65536nodes
Graph500: 17977 GTEPSThe size of graphs

Human Brain Project Graph500 (Huge)

1 trillion
edges

45

Symbolic
Network
Graph500 (Large)

log2(m)40

Graph500 (Medium)1 billion
edges

35

Twitter (tweets/day)
Graph500 (Small)
Graph500 (Mini)

30

Graph500 (Toy)USA-road-d.USA.gr

25

USA-road-d.LKS.gr
20
15

USA-roadd.NY.gr
20

Mobile Phone : SONY SO-01F
Snapdragon S4 1.7GHz 4core: 2GB RAM
1.03GTEPS: 235.06MTEPS/W

USA Road Network
1 trillion
nodes

1 billion
nodes
25

30

log2(n)

35

40

45

No. of nodes

Graph500 “Big Data” Benchmark
Kronecker graph BSP Problem

A: 0.57, B: 0.19
C: 0.19, D: 0.05

November 15, 2010
Graph 500 Takes Aim at a New Kind of HPC
Richard Murphy (Sandia NL => Micron)“ I expect that this ranking may at times look very
different from the TOP500 list. Cloud architectures willalmost certainly dominate a major chunk of part of thelist.”The 8th Graph500 List (June2014): K Computer #1, TSUBAME2 #12
Koji Ueno, Tokyo Institute of Technology/RIKEN AICS

Graph
= Nodes andConnecting
Edges

Reality: Top500 Supercomputers Dominate#1 K Computer
No Cloud IDCs at all#12 TSUBAME2

Elapsed Time (ms)Sparse BYTES: The Graph500 – 2015~2016 – world #1 x 4
K Computer #1 Tokyo Tech[Matsuoka EBD CREST] Univ.
Kyushu [Fujisawa Graph CREST], Riken AICS,88,000
Fujitsunodes,
1500

Communi…
Computati…

1000

73% total exec
time wait incommunication#1 38621.4 GTEPS

500
0660,000 CPU Cores
1.3 Petabyte mem
20GB/s Tofu NW(#7 10.51PF Top500)(Scale 30)65536 nodes

Effective x13
performance c.f.
Linpack≫

64 nodes

List

Rank

GTEPS

November 2013

45524.12(Scale 40)
LLNL-IBM Sequoia
BYTES Rich
TaihuLight
1.6 million CPUs
Implementation
Machine + Superior
10 million CPUs
1.6 Petabyte mem
1.3 Petabyte mem
Top-down only BYTES algoithm

June 2014

1

17977.05

Efficient hybrid

November 2014

2

19585.2

Efficient hybrid

June 2015
June 2018

138621.4

Hybrid + NodeCompression#3 23751 GTEPS #2 23755.7 GTEPS(#4 17.17PF Top500) (#1 93.01PF Top500)

BYTES, not FLOPS!

EBD vs. EBD Calculation (Akiyama Gr, Tokyo Tech)
Highly-sensitive Homology Search Demanded in Metagenomics
Sea

O(m) Reference

Databases

Soil

Human bodyO(n)Sequencer

Meas.
data

~TB

O(m n) calculation
Similarity search>TB

Human oral metagenome analysis:
Input:

Measured Data: 1,284 GB(18 billion DNA reads)
A Reference DB: 15 GB(KEGG GENES DB only)

Output:Search Result: 1,223 GB(Similarity, Ontology)

on a Xeon cluster w/144 cores
0.18 M reads/hour(100,000 hours / 18 billion reads)

x3,178 fasteron K-computer
572 M reads/hour(32 hours / 18 billion reads)

and beyond...

Development of Ultra-fast Homology Search Tools
x100,000 ~ x1,000,000 c.f. high-end BLAST WS (both FLOPS and BYTES)GHOSTZ-GPU

GHOST-MP

Suzuki, et al. Bioinformatics, 2015.

Suzuki, et al. PLOS ONE, 2016.

Kakuta, et al. Int J Mol Sci, 2017

Subsequence sequence clustering

Multithread on GPU

MPI + OpenMP hybrid parallelization

TSUBAME 2.5

100
10
1

BLAST

GHOSTZx 240 faster than conventionalalgorithm

40
20
0
12C+3G

1000

×80-100 faster12C+1G

10000

60

1C+1G

100000__ GHOST-MP80

1C

1000000

TSUBAME 2.5 Thin node GPU
Speed-up ratio for 1 corecomputational time for
10,000 sequences (sec.)(3.9 GB DB、1CPU core)x 70 faster than 1 coreusing 12 cores + 3 GPUs

mpi-BLAST

Retaining strong scaling
up to 100,000 cores

Application Co-design in EBD-CREST
and Challenge for Large-scale HPC on Public Cloud Environment
GHOST-MP system• Original programing by C++ and MPI.
• Original task assignment system (MPIDP) (implemented by Akiyama Gr)
A more general framework is required to facilitate porting andimplementation(1) Using MapReduce w/Matsuoka Gr (EBD CREST Co-design)(2) Using Pwrake/Gfram w/Tatebe Gr (EBD CREST Co-design)(3) Using HPC cloud and container virtualization(2) Using Pwrake/Gfram w/Tatebe Gr
300000(3) Using HPC cloud and container virtualization

8.9x speed up

200000GHOST-MP

150000100000
50000
0

1

2

3

4

Pwrake/Gfarm
5
6
7

node num

1thread8threads

GHOST-MPx speedup

elapsed time [s]250000~80%
strong scaling#virtual machines

Utilization of GHOST in Real-world Projects
Joint researches
1) Tokyo Dental University, Prof. Ishihara“Marker genes exploration of periodontal disease”
2) 7SP Project by KOINOBORI Associates Inc.
“Gut microbiome analysis of mitochondrial diseases”
3) National Agriculture and Food Research Organization (NARO)“Metagenomic analysis of biofilm on electrode in microbial fuel cell”
4) The Institute of Medical Science, The University of Tokyo, Prof. Uematsu“Mouse gut immune system and aging”

Large computing centers5) Kyoto University KEGG
GHOSTX is implemented in “GhostKOALA (software for metagenome analysis)”6) National Institute of Genetics (NIG)
GHOSTX is implemented in “Bacteria genome analysis pipeline DFAST”
7) Nanyang Tech Univ (SCELSE) Singapore & NSCC Singapore
GHOSTZ-GPU has been installed on both computational environment.

18

Utilization of GHOST in Real-world Projects
Tokyo Dental University
Marker genes exploration of periodontal disease (歯周病)
To extract features of healthy/disease states from bacterial taxonomic composition is difficult,because the composition is highly valuable among individuals.

Prof. Ishihara

GHOST series software reduce the computational cost and make metagenome analysis more casually.
Additionally, the diversity of functional gene composition is lower than that of taxonomic composition and this data isapplicable to machine learning method.

Taxonomic annotation

Functional annotation by GHOSTZ-GPU

Summary
Achievements in EBD-CREST Activity
Initial GOAL of Akiyama Gr
1. To develop ultra-fast and high-sensitive homology search tools
2. To generalize the invented scheme, as a common tool for “EBD vs. EBD”
3. To realize of practical life science applications from real-world bigdata
Our ANSWERs
1. GHOST-MP is about x100 faster than mpi-BLAST, and has extreme scaling with its original
efficient load-balancing mechanism (MPIDP). Furthermore, we renovated the engine to employ
further faster GHOSTZ algorithm and implemented GHOSTZ-GPU utilizing multi-GPUs.
2. Architecture co-design to realize common tools for “EBD vs. EBD” has done with Matsuoka Gr.
and Tatebe Gr. In addition, we have also completed the system development to utilize thegeneral-purpose cloud HPC (Microsoft Azure) that evolves in recent years.
3. We have analyzed whole human oral microbiome data on HMP database which has 18 billion
DNA reads (1.3 TB) with high-sensitivity search (published on IJMS journal).
4. Collaborative researches on diseases with multiple research institutions were conducted(UTokyo, TDU, NARO, KOINOBORI Inc.). The GHOST system was installed at multiple research
institutes (Kyoto Univ, NIG, NTU, NSCC).

現代のAIはHPCにより「復興」
2012年6月「キャットペーパー」の衝撃
Googleの「ネコ認識」→機械学習によって、自らネコの概念を獲得し、識別。

実際に獲得された
ネコの「概念」
Le et al.(2012)

ディープニューラルネットワーク

AI研究のブレイクスルー。しかし原理は1970年代からあった

Society5.0の実現に向け 画像認識、ロボット・自た飛躍的な発展
動運転、自動翻訳…etc.

従来の学習法と比べ、(1)学習データ量を増やすとどんどん精度が上
Fig. 2: Andrew Ng (Baidu) “What Data Scientists Should
がるが、(2) その分大量の計算が必要になる
Know about Deep Learning”
なので、計算パワーが100万倍になって初めて可能になった
学習の計算パワーを上げるには、通常のシミュレーションと同じく(1) 個々のCPUでの速度をハード・ソフトの工夫で高速化(2) それを大規模並列化する

スパコンによる
大規模な計算4 Layers of Parallelism in DNN Training• Hyper Parameter Search• Searching optimal network configs & parameters• Parallel search, massive parallelism required• Data Parallelism• Copy the network to compute nodes, feed different batch data,
Inter-Node average => network reduction bound• TOFU: Extremely strong reduction, x6 EDR Infiniband• Model Parallelism (domain decomposition)• Split and parallelize the layer calculations in propagation• Low latency required (bad for GPU) -> strong latency tolerant
cores + low latency TOFU network• Intra-Chip ILP, Vector and other low level Parallelism• Parallelize the convolution operations etc.
• SVE FP16+INT8 vectorization support + extremely high memory
bandwidth w/HBM2
Intra-Node• Post-K could become world’s biggest & fastest platformfor DNN training!

Massive amount of
total parallelism,
only possible via
supercomputing

22

Deep Learning is “All about Scale”
Massive Parallelization is the key•

Data-parallel training with (Asynchronous)
Stochastic Gradient Descent– Replicate network to all the nodes, feed different data, average thegradients periodically– Network All-Reduce Reduction in Megabytes~Gigabytes becomes
the bottleneck at scale– NVIDIA: NVLink Hardware + NICL library (up to 8 GPUs on DGX-1,
16 on DGX-2 w/ NVL Switch)

Fig. 2: Andrew Ng (Baidu) “What Data Scientists Should
Know about Deep Learning”

November 26, 2018

Fig. 3: Simplified DL workflow with ASGD per iteration:
1. Compute gradient
2. Exchange gradients via all-reduce; and3. Update network parameters
Jens Domke

23

Predicting Statistics of Asynchronous SGD Parameters for a Large-Scale
Distributed Deep Learning System on GPU Supercomputers

Background

Proposal• In large-scale Asynchronous Stochastic Gradient Descent • We propose a empirical performance model for an ASGD
(ASGD), mini-batch size and gradient staleness tend to be
deep learning system SPRINT which considers probability
large and unpredictable, which increase the error of trained distribution of mini-batch size and staleness
DNN

Mini-batch size

Objective function E
Mini-batch size-ηΣi ∇Ei

W(t)W(t+1)

4 nodes
8 nodes
16 nodesStaleness=0

Twice asynchronous
updates within
gradient computation

W(t+3)

Staleness

Predicted
Measured

Predicted

W(t+1)
W(t+2)•

Staleness=2-ηΣi ∇Ei
DNN parameters space

Measured(NSubbatch: # of samples per one GPU iteration)

Yosuke Oyama, Akihiro Nomura, Ikuro Sato, Hiroki Nishimura, Yukimasa Tamatsu, and Satoshi Matsuoka, "Predicting Statistics ofAsynchronous SGD Parameters for a Large-Scale Distributed Deep Learning System on GPU Supercomputers", in proceedings of2016 IEEE International Conference on Big Data (IEEE BigData 2016), Washington D.C., Dec. 5-8, 2016

Interconnect Performance as important
as GPU Performance to accelerate DL• ASGD DL system SPRINT (by DENSO IT Lab) and DL speedup prediction
with performance model– Data measured on T2 and KFC(both FDR) fitted to formulas– Allreduce time (∈ TGPU) dep. on#nodes and #DL_parameters

Fig. 4: Oyama et al. “Predicting Statistics of Asynchronous SGD Parameters for aLarge-Scale Distributed Deep Learning System on GPU Supercomputers• Other approaches == similar improvements:– Cuda-Aware CNTK optimizes communication pipeline  15%—23% speedup(Banerjee et al. “Re-designing CNTK Deep Learning Framework on Modern GPU Enabled Clusters”)– Reduced precision (FP[16|8|1]) to minimize msg. size w/ no or minor accuracy loss

TSUBAME 3.0 のシステム概要

全GPU・全ノード・全メモリ階層での
BYTES中心スケーラブルアーキテクチャ

2017年8月本稼働

DDNのストレージシステム（並列FS 15.9PB+ Home 45TB）

フルバイセクションバンド幅のインテル® Omni-Path® 光ネットワーク
432 Terabits/秒 双方向
全インターネット平均通信量の2倍５４０の計算ノード SGI ICE® XA
インテル®Xeon® CPU×２＋NVIDIA Pascal GPU×４
256GBメモリ、２TBのNVMe対応インテル®SSD
47.2 AI-Petaflops, 12.1 Petaflops

TSUBAME3.0 Co-Designed SGI ICE-XA Blade (new)- No exterior cable mess (power, NW, water)- Plan to become a future HPE productLiquid Cooled“Hot Pluggable” ICEXA Blade
Smaller than 1U server,
no cables or pipes

Xeon x 2
PCIe Switch> 20 TeraFlops
DFP
256GByte Memory

PCIe NVMe
Drive Bay x 4

100Gbps x 4

Liquid Cooled NVMe

15 Compute Racks
4 DDN Storage Racks
3 Peripheral & SW
racks
Total 22 Racks

144 GPUs & 72
CPUs/rack
Integrated
100/200Gbps
Fabric Backplane

TSUBAME3.0 became the first large
production petaflops-scale
supercomputer in the world to be #1
on the “Green500” power efficiency W
world ranking of supercomputers14.1 Gigaflops/W is more than x10
more efficient than PCs andSmartphones!

Award Ceremony at
ISC2017 @ Frankfurt
Power Meters

Tokyo Tech / HPE
Benchmarking TeamTSUBAME3: A Massively BYTES Centric Architecture for Converged BD/AI and HPC
Intra-node GPU via NVLink
20~40GB/s

HBM2
64GB
2.5TB/s

DDR4
256GB
150GB/sIntel Optane
1.5TB 12GB/s(planned)
NVMe Flash2TB 3GB/s

Terabit class network/node
800Gbps (400+400)
full bisection

Intra-node GPU via NVLink
20~40GB/s

Inter-node GPU via OmniPath
12.5GB/s fully switched

Any “Big” Data in thesystem can be moved
to anywhere via
RDMA speeds
minimum
12.5GBytes/s
also with Stream 16GB/s PCIe
16GB/s PCIe
Processing
Fully Switched
Fully Switched
Scalable to all 2160
GPUs, not just 8~4 Terabytes/node Hierarchical Memory for Big Data / AI (c.f. K-compuer 16GB/node) Over 2 Petabytes in TSUBAME3, Can be moved at 54 Terabyte/s or 1.7 Zetabytes / year31TSUBAME3: A Massively BYTES Centric Architecture for Converged BD/AI and HPC
Intra-node GPU via NVLink
20~40GB/s

HBM2
64GB
2.5TB/s

DDR4
256GB
150GB/sIntel Optane
1.5TB 12GB/s(planned)
NVMe Flash2TB 3GB/s

Intra-node GPU via NVLink
20~40GB/s

Inter-node GPU via OmniPath
12.5GB/s fully switched

Any “Big” Data in thesystem can be moved
to anywhere via
RDMA speeds
minimum
12.5GBytes/s
also with Stream 16GB/s PCIe
16GB/s PCIe
Processing
Fully Switched
Fully Switched
Scalable to all 2160
GPUs, not just 8~4 Terabytes/node Hierarchical Memory for Big Data / AI (c.f. K-compuer 16GB/node) Over 2 Petabytes in TSUBAME3, Can be moved at 54 Terabyte/s or 1.7 Zetabytes / year

32

33

DRAGON: Breaking GPU Memory Capacity
Limits with Direct NVM Access*
Pak MarkthubT, Mehmet E. BelviranliO, Seyong LeeO,
Jeffrey S. VetterO, and Satoshi MatusokaR,T
TTokyo Institute of Technology
OOak Ridge National Laboratory
RRIKEN Center for Computational Science*Full paper (same title) will be published at SC18
Code is available at https://github.com/pakmarkthub/dragon

Visit us at the poster session

Motivation & Solution• Motivation

•
•
•
•

34
Require Pascal GPU or newer
Not require NVLINKWork on x86 and other CPU architectures
Support remote file storage• Problem sizes grow larger than GPU and host memory• Need complex GPU algorithms to handle out-of-core, data partitioning…
• NVM: Non-volatile, large capacity, high bytes/$, high IO BW• Solution  DRAGON: Direct Resource Access for GPU over NVM• Enable “mmap” to from files on NVM to GPU via “dragon_map()”• Merge GPU memory, host memory, and NVM spaces with strong and weak data coherency• On-demand paging with stream-like prefetching to improve data transfer performance• Key Contributions:• Transparently extends GPU memory space to cover NVMs (no need to modify GPU kernels)• Eliminate manual buffer management on data larger than the system memory• Significantly improve UM performance up to 2.3x

35

DRAGON: Toward Unlimited GPU Memory
Highlights
Enable mmap from files to UM space
GPU can access data larger than the host memory
UM coherency between GPU, CPU, and NVMe
Modified nvidia-uvm submodule to cover until storage
libdragon

Map & access NVMe

libcudart

ioctl

nvidia-uvm

dragon-driv

Driver

••••Ex. Code
dragon_map("data.mem", size, D_F_READ, &data);
dragon_map("result.mem", size, D_F_WRITE, &result);…
compute_gpu <<< blocks, threads >>> (data, result);
compute_cpu(data, result);

vfs
blockdevGPU

NVMe Dev

nvme

NVM  DRAM

Page fault fromGPU load/store

DRAM  GPU Mem

Data streaming with 2-level prefetchingResults: pathfinder (dynamic proc)
DRAGON is the only existing
solution that enables out-of-coreprocessing with user-oblivious
addressing

Code is available at:https://github.com/pakmar
kthub/dragon

BetterRequire no modification toexisting CUDA kernels in order
to use DRAGONNormalized w.r.t. Unified Mem

2.3x

Overlapping of computation anddata transfer is implicitly
provided by DRAGON
regardless of GPU algorithms#1: cudaMemcpy36#2: hostreg#3: Unified Mem#4: DRAGON

(video action recognition; 3D Conv; batch size 1; ran 30 batches)

Case Study: C3D-UCF101Net (Caffe)
Host mem

Enable complex 37
algorithms that requirelarge workspace to run

Out-of-core

Better

3.5x faster than using
CPUs for out-of-core
GPU mem

DRAGON shows unnoticeable
overhead w.r.t. original Caffe

Normalized
w.r.t. DRAGONDistributed Second-order Optimization Using
Kronecker-factored Approximate Curvature
for Deep Convolutional Neural Networks
Yohei Tsuji1, 2, Kazuki Osawa1, Yuichiro Ueno1,
Rio Yokota1, Satoshi Matsuoka3, 1
1.

Tokyo Institute of Technology

2.

AIST-Tokyo Tech Real World Big-Data Computation Open Innovation Laboratory,
National Institute of Advanced Industrial Science and Technology

3.

RIKEN Center for Computational Science

In progress work, partially presented at “情報処理学会 第165回 HPC研究会（SWoPP2018）”

Motivation• Demand of large scale distributed deep learning• Accuracy degradation on large batch-size training• Ad hoc techniques in SGD to avoid the degradation• More mathematically rigorous method: K-FAC• Approximated second-order optimization method• Potential to avoid the degradation more generally• Requires larger computational time compared to SGD →motivation of using parallel computation

Need an implementation of K-FAC which
can scale out on 1000 or more GPUs

Training ImageNet in Minutes
RioYokota,Kazuki Osawa,YoheiTsuji,Yuichiro Ueno,Hiroki Naganuma,Shun Iwase,Kaku Linsho

Tokyo Institute ofTechnology+Akira Naruse (NVIDIA)

GPU/KNLTSUBAME3.0

ABCI

Source Ben-nun & Hoefler https://arxiv.org/pdf/1802.09941.pdf

Facebook
Preferred Networks
UC Berkeley
Tencent
Tokyo Tech + AIST

512
1024
2048
2048
4096

Time
30 min15 min14 min
6.6 min? minParallel Scalability
Reduce-Scatter for each layer

Cross-over from data-parallel to model-parallel

Kronecker Factorization
Cross-over from model-parallel to data-parallel

Inter-node RabenseifnerAllReduce
Intra-node RingAllReduce
PCH

SSD
DIMM
DIMM
DIMM
DIMM

x16 PCIex4 PCIe

PLX

x16 PCIe

CPU 0

QPI
DIMM
DIMM
DIMM
DIMM

OPA HFI

x16 PCIe

DMI

x16 PCIe

OPA HFI

x16 PCIe

GPU 1

GPU 0
NVLink

CPU 1

GPU 3

GPU 2

x16 PCIex16 PCIex16 PCIe

PLX

x16 PCIex16 PCIe

OPA HFIOPA HFI

Imagenet1k,ResNet50
Local batch size:32

• Using data-parallel AND model-parallel• Reducing computation time with Tensor Core,
half-precision communication in NCCL, …

Scaling out experiments from
1 GPU to 4096 GPUs, almost
linear

Accuracy after 15 epochs on
ImageNet ILSVRC2012 dataset,
BS=8192, top-1 validation
ACC=69.14%

data-parallel• Implementation on Chainer+ChainerMN

model-parallel

Implementation and Results

Averaged second per iteration,achieved 23.3% speedup using
Tensor Core and half-precisioncommunication in NCCL

Center and rightfigures are takenfrom 1024 GPUs
executions

Sekijima Lab., Tokyo Institute of Technology

Drug Discovery Process w/AI on TSUBAME3[M. Sekijima, Tokyo Tech.]
Hit discovery

Process of drug discovery

Optimization
4.5■: Focus

Preclinical

Clinical trials

Approvalin this study

1.0

6.5

1.5

Time [year]

Total time: 13.5 years1
Total investment: 2.6 billion USD2

Aim to reduce 30% of costby Computer-aided drug design

Method: docking-based virtual screening

Problem: prediction accuracy

Identify hit compounds from large (<10M) chemical libraryby simulating/evaluating protein-compound interaction＋→

Decrease false positive through recognizing
interaction pattern by machine learning
Almost the same docking score

Binding energy: -9.52

Scoring function(Glide3))
Active compounds

Inactive compounds

1) Steven MP et al. (2010) 2) Asher M (2014) 3) Richard AF et al. (2004)

Sekijima Lab., Tokyo Institute of Technology

SIEVE-Score: method and evaluation• Extract interaction energy vector from docking (Glide was used)• hydrogen bonding, electrostatic and van der Waals energy

between ligand and each residue of target protein• Train random forest with known experiments,• Re-rank by predicted probability of being “active”

Evaluation: DUD-E dataset4)• 102 proteins, including cancer or hypertension targets• 40,490 active, 1,414,640 inactive compounds in total• 1,455,130 docked structures are learned• 5-fold cross validation for each protein

TSUBAME2.5&3• Metrics: enrichment factor 1%

The ratio of identified active compounds in top 1% compared to random prediction

Flow of SIEVE-Score

4) Mysinger MM et al (2012)

Sekijima Lab., Tokyo Institute of Technology

SIEVE-Score outperformed other methods• SIEVE-Score improved most of

the results against Docking
for all target in DUD-E (right).
• In the EF 1%, SIEVE-Score is 42.1
against Docking's 10.3.
• hit compounds can be found 4 times more efficiently.

SIEVE-Score Mean: AUC 0.89
SIVE-Score vs Docking
Docking: AUC 0.638

Target: serine/threonine-protein kinase AKT (cancer target)日本のビッグデータ・AIのエコシステムの現状
最先端のアルゴリズム・マシンインフラ・データが揃う事が必須だが、インフラ部分が一番欠けている。
Joint
RWBC
Open Innov.
Lab (OIL)

AI ベンチャー
AIST-AIRC(Director: Matsuoka)

AIやビッグデータの

Riken
-AIP

AIアルゴリズム
やソフトの
研究開発

企業におけるAI・機械学習に関する研究開発と応用
企業でのデータ活
用の研究

NICT- 研究所や大学とのUCRI 連携

大量のデータが企業や研究所にあるが多くが捨てられている
自動運転

B

Massive Rise in Computing
Requirements (1 AI-PF/person?)

ＨＰＣではスパコンという公共インフラおよびその運用体制が確立しているが、ビッグデータ・AIでは欠落

工場・ロボット

流通・リテール

AIおよびデータ 凄まじいデータと
計算の要求
のインフラ

ビッグデータ通信・IoT

オープンAIインフラストラクチャ
AI橋渡しクラウド（ABCI）の構築
https://abci.ai/

国立研究開発法人 産業技術総合研究所
人工知能研究センター (AIRC)

49

世界最大級・超省電力・オープン
ＡＩインフラストラクチャ
AI Bridging Cloud Infrastructure

世界トップレベルの計算処理能力とデータ処理能力

ＡＩとビッグデータのアルゴリズム、ソフトウェア、
応用開発のためのオープンかつ専用の計算インフラストラクチャ

我が国における産学官共同のＡＩ研究開発を加速するオープンイノベーションプラットフォーム

東京大学 柏IIキャンパス

理論ピーク性能: 550 PFlops (半精度)
37 PFlops (倍精度)
実効性能(Linpack) : 19.88 PFlops (世界5位, 国内1位)
12.054 GFlops/W (世界8位)
使用電力(最大): 2.3 MW
年間平均PUE: 1.1以下 (推定値)
2018年8月1日運用開始

50

ABCI: 強みを伸ばし、いち早い導入を試す場とモデルの提供
AIへの関心と導入状況：「既に導入している」はまだ3.1％。「実証試験を行っている」7.1％を合わせて１割程度。

AIを試す
場がない
（出典）情報処理推進機構「AI社会実装推進調査報告書」（2018年6月19日）より作成

AIの実需要がないJanuary 16, 201751

October 30, 201752

パネマジ化

53

54

55

57

58

59

60

ＡＩデータセンター棟“1ラック70kW級スパコンのコモディティ化を可能に”
クーリングポッド

屋外設備置場
高圧受電設備 (3.25MW)

サーバ室: 19m x 24m x 6m

 低コストで軽量な建物

UPS
将来の拡張用 耐荷重2トン/m2のコンクリートスラブ上にクーリングポッドとラックを直接設置

18ラック（UPSバックアップ）

アクティブ
チラー

パッシブ
冷却塔

冷却能力:
200kW

冷却能力:
3MW

72ラック

 ラック数• 初期: 90 (ABCIは41台使用)• 最大: 144

 電力容量• 3.25 MW (ABCIは最大2.3MW使用)

将来の拡張用 冷却能力• 3.2 MW (夏季の最小値)• ラックあたり水冷60kW＋空冷10kW
Data Center Image従前型のデータセンターとTSUBAME3およびABCI の比較
AI 向けのIDCの効率を100倍以上改善
性能 > 400~600
効率 > 200~300(Pascal => VoltaへのUGを仮定)

従来型のクラウドIDC~10KW/rack PUE 1.5~2
15~20 1U Xeon Servers
2 Tera AI-FLOPS(SFP) / server
30~40 Tera AI-FLOP / rack

T3 / ABCI IDC~60KW/ラック PUE 1.1以下~36 T3 evolution servers in ABCI
~500 Tera AI-FLOPS(HFP) / server~17 Peta AI-FLOPs / rack62

ポスト「京」 ～The Game Changer～

1. コ・デザインによる京から受け継ぐシミュレーションでの優位性◎アプリにおける高演算性能：京と比較で最大100倍以上の性能向上◎重点・萌芽課題における多くの科学的成果創出の準備・期待◎高性能の達成と、容易なプログラミング・高実用性の両立

2. ポスト京の新たな技術イノベーション◎チップ自身の高演算性能（高速メモリ→高メモリバンド幅）
CPUチップ単位で、多くのHPC＆Society5.0アプリで従来CPU数倍の性能◎高い省電力・グリーン性能

マシン自身だけでなく半導体やITのテクノロジー
でも世界をリード「パワーノブ」含む省電力技術導入に、汎用CPUでは世界トップレベルの高効率◎Arm「エコシステム」の充実

年間30億個生産されるArmプロセッサの命令セットを採用
クラウドにも展開が容易、今回成果のSVEのグローバルスタンダード化◎Society5.0アプリへの展開

ビッグデータ、AI、CAE/EDA、次世代コンピュータセキュリティ・ブロックチェーンなどのアプリにおいて、GPU並みからそれ以上の性能を発揮すると期待。

ポスト「京」は“二刀流”

ARM：スマホ等の機器、車載チップ等で大きな実績

ビッグデータ・AIへの展開

更にインフラ、および技術としても世界に展開普及もCPU
For thePost-K

supercomputer63ポスト京における新たな技術イノベーション：サマリ

1. 高速メモリテクノロジによる高メモリバンド幅、それに伴うHPCアプリの高性能演算性能


コ・デザインアプリに加え、世界のベンチマーク調査から、多くのHPCアプリはメモリ性能律速、一部は演算性能律速→高いメモリ性能と、それとバランスした演算性能の向上が重要

ポスト京CPUと現存トップサーバCPUと比較：何倍ものメモリ性能と、比較優位な演算性能

超高速メモリに対応し(CPUでは世界初) 、演算性能も上げたCPUの設計が成功のカギ

2. 高い省電力・グリーン性能

電力最適化設計、クロックゲーティング、パワーノブ等の高性能省電力設計



CPUでは世界トップレベルの電力性能比、及び高効率冷却→クラウドなど一般ITにも適用可

3. ARM サーバ&HPC エコシステムの充実→使いやすさ・広範囲な成果普及

年間のプロセッサの生産量：x86 3億個 vs. ARM 30億個



高性能エコシステムの急速な確立 (e.g. Cavium, HPE, Sandia, Bristol, CEA,…)

ポスト京 SVE(Scalable Vector Extension) ->スパコン・クラウドからIoTまでグローバル標準

3. Society5.0：ポスト「京」の性能優位性AI・機械学習の事例

次世代の深層学習によるAIは莫大な演算量→大規模なHPCが必要
Post-Kプロセッサ：FP16などのAI専用命令と、高バンド幅メモリ =>深層学習の中心である「畳み込み演算」の優れた性能と、大規模な並列化のスケーラビリティ
世界トップのAIマシンの可能性、Blockchain Securityなど他のScociety5.0アプリでも同様

64

ポスト「京」 におけるコ・デザイン

システムとアプリケーションの協調的な開発
手法（Co-design）「計算の科学」「計算による科学」
・重点9課題：アプリケーション対象として健康長寿、
防災・減災、エネルギー、ものづくり分野等の様々な社会的・科学的課題を選定

CPU
For thePost-K

supercomputer

各種アプリケーション 各アプリ特性を加味したの様々な計算上の特性 システム設計、及びそれに適したアプリの最適化
を一般化・特徴化

 理研CCS(旧AICS)と開発企業、及び重点課題機関の密な研究開発体制

 コ・デザインにより多くのアプリから9つの代表的な「ターゲットアプリケーション」の選定

 複数のターゲットアプリで、コ・デザインの結果、京と比較して100倍以上の性能向上
 プログラミングを含む使い易さの向上・高い電力性能などの技術成果も

現在のアーキテクチャにおけるアプリの性質
米国ECP (Exascale Computing Program)の Proxyapps+Post-K miniappsの解析結果
Broadwell Xeon

Compute Performance (GFlop/s)

128

CANDLEmVMC-MINI

32
16
8
4

MINI
SW4liteNTChemMINI
NekboneBenchmarks

64

XSBenchAMG

MODYLAS-

CCS QCDminiFE

CoMD

FFB-MININICAM-DCMINI

SWFFT

LaghosminiAMR2
1
0.0625

0.25

1

4

16

64

FLOPS/BYTES ratio66

1. 高いメモリバンド幅によるHPCアプリの高演算性能
 種々のベンチマークから、多くのHPCアプリケーションはメモリ律速

ポスト京：FIBREベンチマーク、DoE ECP ベンチマーク
 メモリコントローラのビジー率の高さ ハイエンドXeon比較でポスト京プロセッサが圧倒的メモリバンド幅 Xeon: 6~8チャンネル DDR4、~100GB/s (STREAM)
 ポスト京: 4 Stack HBM2 2Ghz, 1TB/s(スペック),~840GB/s (STREAM)
 CPUとして世界初の採用、バンド幅・レーテンシ・コヒーレンシをすべて満たす世界トップクラスのメモリコントローラの設計
 最新・ハイエンドのIntel Xeon Platinum 比で、計算で2倍、メ
モリでは8倍の性能、非常に省電力 Volta GPUクラスの性能や電力性能比
 メモリ周りの最適化：単純にHBMをボルトオンしたのではない 大規模HPCだけでなく、ビッグデータ・AI・セキュリティなどSociety
5.0 アプリで高性能

部外秘

PostK Proc: HBM2
256GB/s x 4Skylake Xeon
6ch DDR4 2666
21.3GB/s x 6 / chip67

A64fx Performance - A0 Silicon A64FX boosts performance up by microarchitectural enhancements,
512-bit wide SIMD, HBM2 and process technology• > 2.5x faster in HPC/AI benchmarks than SPARC64 XIfx (Fujitsu’s previous HPC CPU)
• The results are based on the Fujitsu compiler optimized for our microarchitecture and SVE
A64FX Benchmark Kernel Performance (Preliminary results)
8

HPC

Throughput

AI

Application
Kernel(DGEMM / Stream)

9.4x(PRIMEHPC FX100)

Normalized to
SPARC64 XIfx

6
Memory B/W

512-bit SIMD

Combined
Gather

L1 $ B/W

L2$ B/W

INT8
dot product

4

2

2.5TF

830

3.0x

3.4x
2.8x

2.5x

GB/s

(>90%)(>80%)

DGEMM

Stream
Triad

0
Fluid
dynamics

Atomosphere

Seismic wave
propagation

Convolution
FP32

Convolution
Low Precision(Estimated)

Baseline: SPARC64 XIfx ( PRIMEHPC FX100)
14

All Rights Reserved. Copyright © FUJITSU LIMITED 2018

Post-K A64fx A0 (ES) performance
Performance / CPU

Machine Performance (HPC)

Peak TF(DFP)

Peak Mem.
BW

StreamTriad

TheoreticalB/F

Post-K A64fx(A0 Eng.
Sample)2.764/
3.072

1024GB/s

840GB/s

0.37/
0.33

94 %

Intel KNL

3.0464600GB/s

490GB/s

0.20Intel Skylake

1.6128127.8GB/s

97 GB/s

0.08

NVIDIA V100
DGX-1

7.8

900 GB/s

855GB/s

0.12

GF/W

Network BW
Per Chip87.7 %

GPU
level

40.8GB/s(6.8x
6)

66%

54.4 %

4.9

12.5 GB/s

80 %66.7 %

4.5

6.2GB/s

58 % 15.113

150GB/s
6.2GB/s

DGEMM
Linpack
Efficiency Efficiency

Post-K Chassis, PCB (w/DLC), and CPU Package
230 mmCPU
For the280mm

Post-K

60mm

supercomputer

60mm

W 800㎜
D1400㎜
H2000㎜
384 nodes
FUJITSU CONFIDENTIAL

CPU Package
CMU

A0 Chip Booted in June
Undergoing Tests
Copyright 2018 FUJITSU LIMITEDAn Overview of Post-K Hardware
 Compute Node, Compute + I/O Nodeconnected by 6D mesh/torus Interconnect

 3-level hierarchical storage system


1st LayerCache for global file systemTemporary file systems- Local file system for compute node
- Shared file system for a job

2nd Layer

3rd LayerLustre-based global file system
Storage for archive>100,000 nodes
Approaching 10 million
CPU cores

20018/6/26

71

2. ARM サーバ&HPC エコシステムの充実

 年間のプロセッサの生産量：x86 3億個 vs. ARM 30億個
 サーバチップのハードウェアエコシステムの確立
Cavium: 2018年5月に製品版の量産を発表Competition: Intelに対するSecondary ChoiceSustainability: ARM ベンダーのSecondary choice

サーバベンダーのコミット：HPE, Cray, Fujitsu

メジャークラウドのコミット：Microsoft, Google, etc.

 HPCソフトウェアエコシステムの確立

今回新たにSVE (Scalable
Vector Extension)を世界初めて提案・実装。
これはARM v8の正式プ
ロファイルであり、v9では組み込みなどへも普及
本プロジェクトの成果がスパコン・クラウドから
IoTまでグローバルデ
ファクトスタンダードへ

米国：DoE Sandia/Los Alamos NL, NERSC



欧州：European Exascale, CEA(仏), BSC(西), EPCC&Bristol (英)

中国：NUDT-Tianhe 3、上海交通大学
72

ARM HPC ECOSYSTEM事例: 英国(slides courtesy Prof. Simon Mcintosh-Smith) 英国はARMの本拠地、採用に積極的

 多くのアプリやソフトが移植中

 その他フランス(CEA), スペイン
(BSC), Sandia
NL(US), など

 ほぼ全てのx86系の

ソフトはそのまま動作OS, 言語、ツール

複数のコンパイラ

73

Society5.0：ポスト「京」の性能優位性
AI・機械学習の事例

高性能な深層学習の畳み込み演算
今後の大規模なデータ学習を必要と◆高性能低精度演算(FP16, INT8)するAI研究において、世界トップの◆高メモリバンド幅
マシンとなる！
◆高性能・スケーラブルネット
ワーク① チップ単位の高い② 高いネットワーク通信性能による畳み込み演算性能
超並列化：Data+Model Parallelism
FP16+INT8 と高速メモリ
CPU
For thePost-K

supercomputerCPU
For thePost-K

supercomputer

畳み込み演算の適切な選択（FFT, Winograd）+高メモリバンド幅＋低精度演算機能により、GPUに匹敵or超える性能が見込める。

CPU

CPU

Post-K

Post-KFor thesupercomputer

For thesupercomputer

TOFU超高速ネットワーク
スケーラブルかつ低レーテンシ通信性能によって、大規74
模な並列化(model & data)が見込めるAIにおけるポスト「京」の性能優位の可能性
ポスト「京」CPU
チップ
の性能
通信
スケー
ラビリ
ティ

ソフト
ウェア
エコシ
ステム

FP16+INT8、および高メモリバンド幅⇒高い畳み込み演算性能 ◎通信のスケーラビリティ◎(数十万)

6-D Torus, 高インジェクションBW
モデル、データ並列とも大規模にスケール

エコシステム：ARM ◎ほとんどのXeon向けのフレームワークは容易に移植可能(ただし性能を得るには研究開発が必要)

演算性能

TPU

GPU

※Google オリジナルのプ
ロセッサ◎

演算性能 ◎
ただしproprietaryで入手不可能

ただしトップエンドのGPUは大変高額△ (up to 16 on DGX2)
16を超えると大幅にダウン、かつModel 並列には適さないエコシステム ◎(ただしNVIDIAのみ、またAI以外では〇)？？ (2D Proprietary
Torus)、京ほどはス
ケールしない× Google独自開発
TensorFlowのみ対象
他のアプリ領域の柔軟性が低く、開発したプログラムの他サービスへの移行困難

75

Why memory and NW BW is important for scalable learning?

• メモリバンド幅

アルゴリズム (cuDNN 7.0.1, P100-SXM2)

畳み込みアルゴリズムの演算量の比較(GEMM)

Oyama, Yosuke, et al. ”μcuDNN: Accelerating Deep Learning Frameworks
with Micro-Batching." arXiv preprint arXiv:1804.04806 (2018).

積和回数:
Winograd
∧
FFT
∧
GEMM

Tal Ben-Nun and Torsten Hoefler. "Demystifying Parallel andDistributed Deep Learning: An In-Depth Concurrency
Analysis." arXiv preprint arXiv:1802.09941 (2018).

FFTによる畳み込み
Mathieu, Michael, Mikael Henaff, and Yann LeCun. "Fast training ofconvolutional networks through ffts." arXiv preprint arXiv:1312.5851 (2013).

メモリ要求:

実効性能:

1x1

3x3

フィルタサイズ5x5〜

im2col (GEMM)による畳み込み
Chetlur, Sharan, et al. "cudnn: Efficient primitives fordeep learning." arXiv preprint arXiv:1410.0759 (2014).

Winograd
∧
GEMM
∧
FFT

Winogradのアルゴリズム

Liu, Xingyu, et al. "Efficient sparse-winograd
convolutional neural networks." arXiv preprint
arXiv:1802.06367 (2018).

Validation loss

多くのメモリアクセスを伴うFFTによる畳み込みと高バンド幅メモリを組み合わせることで、CNNのさらなる高速化が可能とな データ一時保存用のメモリ量 [MiB]
AlexNetで使用される畳み込み演算ることが予想される。

訓練データサイズ•

NVIDIA GPU上での畳み込み演算を行うデファクトスタンダードであるcuDNNライブラリでは、多くの畳み込みレイヤーについてFFTを使った畳み込みがWinogradのアルゴリズムと遜色のな
い計算性能となることがわかっている。

• ネットワークバンド幅• 深層ニューラルネットワークの精度を線形に向上させるためにはネットワークの規模と訓練データをそれぞれ指数関数的に増加させる必要があることが分かっている。

計算時間 [ms]•FFT

畳み込み演算は畳み込みニューラルネットワーク (CNN) の根幹をなす基礎的な演算であり、学習/推論の計算時間の大部分を占めると言われている。これを高速に行うためにFFTを用いた畳み込みやWinogradのアルゴリズムと呼ばれる少ない積和演算で畳み込みを実現するアルゴリズムが用いられている。

Better•Winogradネットワーク規模
ネットワーク規模
LSTMを用いた言語モデルのネットワーク規模と精度
Hestness, Joel, et al. "Deep Learning Scaling is Predictable, Empirically." arXiv preprint
arXiv:1712.00409 (2017).

• これは超大規模並列計算機を用いないと、より高い精度の学習ができなくなる日がくることを意味している。今後、自動運転や医療など国際競争力の基盤となる人工知能技術で必要となる
深層学習は従来の学習よりも遥かに難しい問題を解くことになり、それには超大規模並列計算機が必ず必要になる。
• しかし、このような大規模分散並列計算機上で深層学習を行う場合、データ並列、モデル並列、
非同期最適化を組み合わせたとしても必ず膨大な通信がノード間で発生することになる。現時点で既に通信がボトルネックとなっており、GPUなどの演算性能が優れていても通信が速くなければ短時間で学習することは困難である。
今後この傾向はますます強くなると予想される。

どのようにして適切な畳み込みカーネルを選択するか
A64fxにおける深層学習の超高速化
 ETHとの共同研究：Micro Batching [Oyama, Tan,

Hoefler & Matsuoka, IEEE Cluster2018]










Evaluation: WR using Dynamic Programming• μ-cuDNN achieved 2.33x speedup on forward convolution of
AlexNet conv2

DNNにおいて、micro-batch技法を用い適切な畳み込みカーネルを自動選択
速度とメモリのバランスを最適化
GPU上の実験では、多くの場合GEMMでなくWinograd or FFTが選ばれる

cudnnConvolutionForward of AlexNet conv2 on NVIDIA Tesla P100-SXM2Workspace size of 64 MiB, mini-batch size of 256
Numbers on each rectangles represent micro-batch sizes

Evaluation: WD using Integer LP

cuDNNのラッパーなので、全てのDNNフレームワークに透過的に適用可能
ポスト京チップでは、(1)更にWinograd/FFTの選択のケースが増える、(2)その場合GPUと速度は同じ、GEMMを盲目的に用いている場合よりは高速化

A desirable configuration set of AlexNet conv2 (Forward)
Mini-batch size of 256, P100-SXM2
Each bar represents proportion of micro-batch sizes and algorithms

77

Large Scale simulation and AI coming together[Ichimura et. al. Univ. of Tokyo, IEEE/ACM SC17 Best Poster]

130 billion freedom
earthquake of entire Tokyo
on K-Computer (ACM
Gordon Bell Prize Finalist,
SC16,17 Best Poster)Soft Soil

Earthquake

？

<100m
Candidate
Underground
Structure 1

Too Many Instances

AI Trained by Simulation
to generate candidate
soft soil structure

Candidate
Underground
Structure 278

我が国の AI インフラの拡充に向けて
東工大・産総研・理研によるビッグデータ・AI の HPC加速
設計中
ここでも革新・実用・継続が重要
2017年8月

運用中

x5.8(産総研AIRC・NEC)8.2 AI-ペタフロップス

運用中
2015年10月

TSUBAME-KFC/DL(東工大・NEC・NVIDIA)
1.4 AI-ペタフロップス

2017年4月
RAIDEN(理研AIP・富士通)
4.1 AI-ペタフロップス

ABCI (産総研AIRC)

2020年代2018年6月

x2.8~4.2

x5.8 47.2 AI-PF (Tsubame2.5合算だとAIST AI クラウド 65.8 AI-PF)
2017年4月

運用中

TSUBAME3.0(東工大・HPE)

x5.0~7.7

2019前半“ExaAI”

調達・構築中運用中

550 AI-ペタフロップス
以上(実際は数百?)

ポスト京も世界
トップクラス
AI-エクサフロップ
ス性能へ

IDCや設備もＡＩ用に革新2016年10月 ホワイトハウス AI戦略白書「AIにおける最先端のアルゴリズム・ソフ
ト・ハードさらにそれらのコデザインのためにはオープンなインフラを中心としたコミュニティ形成が必須」

An Overview of System Software Stack
Easy of use is one of our KPIs (Key Performance Indicators)
Providing wide range ofapplications/tools/libraries/compilers

Linux Distribution
Eco-System

Fortran, C/C++, OpenMP, Java, …

Batch Job System

Math libraries

Hierarchical File System

Tuning and Debugging Tools

Parallel File System

Parallel Programming Environments
Communication Application-oriented
XMP, FDPS, …
MPI
File I/O
Process/Thread
PIP

Low Level Communication

File I/O for
Hierarchical Storage
LLIO

Multi-Kernel System: Linux and light-weight kernel (McKernel)
Armv8 + SVE
20018/6/26

80

Post-K Programming EnvironmentPrograming Languages and Compilers
provided by Fujitsu

Fortran2008 & Fortran2018 subset



C11 & GNU and Clang extensions



OpenMP 4.5 & OpenMP 5.0 subset



Java



Parallel Programming Language &Domain Specific Library provided by
RIKEN





Process/Thread Library provided by RIKEN


20018/6/26

E.g., Python+NumPy, SciPy

Communication Libraries


MPI 3.1 & MPI4.0 subset


pnetCDF, DTF, FTAR

Math Libraries
uTofu (Fujitsu), LLC(RIKEN）

File I/O Libraries provided by RIKEN
Open MPI base (Fujitsu), MPICH (RIKEN）

Low-level Communication Libraries


XcalableMP
FDPS (Framework for Developing Particle
Simulator)

Script Languages provided by Linux
distributor

C++14 & C++17 subset and GNU and
Clang extensions

GCC and LLVM will be also availableが運用する働している。

BLAS, LAPACK, ScaLAPACK, SSL II（Fujitsu）
EigenEXA, Batched BLAS （RIKEN）

Programming Tools provided by FujitsuProfiler, Debugger, GUI

PiP (Process in Process)81