GPU・FPGA複合計算による次世代HPC/AI加速技術
朴泰祐
筑波大学・計算科学研究センター
taisuke@ccs.tsukuba.ac.jp（共同研究：小林諒平・藤田典久・山口佳樹・梅村雅之・吉川耕司）

2022/05/25

1

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

HPCにおけるアクセラレータ利用
n 近年のHPC (High Performance Computing) 分野ではアクセラレータの利用が増えている−Top500(2021/6)おける上位10マシン中7マシンが搭載している−特にGPUの利用が主流
n 主流であるGPUの特徴
GPU

FPGA

理論性能
(FLOPS)

◎

○

並列性

空間
並列

パイプライン
並列

ノード間
通信性能

△

◎

−高い並列演算性能とメモリバンド幅をもつ−しかし並列性が十分でないと性能がでない
条件分岐，並列性の不足，ノード間通信
n そこでFPGAの利用−アプリケーションに特化した回路
パイプライン並列−ハイエンドなFPGAは高速なノード間通信性能をもつ−高位合成環境が発達

22022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

GPU vs FPGA as HPC solutions
deviceGPU

FPGAparallelization

SIMD (x multi-group)

pipeline (x multi-group)

standard FLOPS😃😃 (1000x cores)😃 (~100x pipeline)

conditional branch😢 (warp divergence)😃 (both direction)

memory😃😃 (HBM2e)

😢 (DDR)→😃 (HBM2)interconnect

😢 (via host facility)😃😃 (own optical links)

programming😃 (CUDA, OpenACC, OpenMP)😢 (HDL)→😏 (HLS)

self-controllability😢 (slave device of host CPU)😃 (autonomic)

HPC applications

😃 (various fields)😢 (not much)3

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of TsukubaGPU・FPGAヘテロ演算加速プログラム
n GPUとFPGAを相補的に用いることで高速

化をめざす

CHARM

−CHARM (Cooperative Heterogeneous
Acceleration with Reconfigurable
Multidevices)−それぞれの強みを最大限利用
strong scalingにおけるボトルネックを解消−マルチフィジックスシミュレーション
複数の物理現象の相互作用を考慮したシミュレー
ションであり，様々な特性の演算が出現

+

FPGA

n プログラミング手法が大きな課題−高レベルFPGA記述−GPUとFPGAの協調計算を自然に記述−FPGAコンパイラが鍵GPU

42022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

CHARM: Cooperative Heterogeneous Acceleration with
Reconfigurable Multi-devices
Basic cluster with GPUs (by InfiniBand)

multi-physics/multi-scalecomplicated problem

invoke GPU/FPGA kernsls
data transfer via PCIe
(invoked from FPGA)
GPU
CPU
FPGA
comp.

PCIe

comm.

CPU

PCIe
Application oriented
FPGA-FPGA communication

GPU
FPGA
comp.
comm.

100Gbps direct optical link

Cooperative computing with GPU and FPGA
5

2022/05/25

PCクラスタOSSワークショップ（FPGA）

FPGA Network

Center for Computational Sciences, Univ. of Tsukuba

Cygnus: world first multi-hybrid cluster with GPU+FPGA@ CCS, Univ. of Tsukuba (deployed by NEC)

Cygnus supercomputer at Center for Computational Sciences, Univ. of Tsukuba (Apr. 2019~)
85 nodes in total including 32 “Albireo” nodes with GPU+FPGA (other “Deneb” nodes have GPU only)
6

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

Single node configuration (Albireo)• Each node is equipped withboth IB EDR and FPGA-direct
network• Some nodes are equippedwith both FPGAs and GPUs,
and other nodes are with
GPUs only

SINGLENODE(with FPGA)CPU

Network switch(100Gbps x2)

HCA

2022/05/25

PCクラスタOSSワークショップ（FPGA）

CPU

HCA

HCA

HCA

PCIe network (switch)

PCIe network (switch)

G
P
U

G
P
U

G
P
U

FPGA

Inter-FPGA
direct network(100Gbps x4)

7

Network switch(100Gbps x2)

G
P
U

FPGA

Inter-FPGA
direct network(100Gbps x4)

Center for Computational Sciences, Univ. of Tsukuba

Single node configuration (Deneb)• Each node is equipped withboth IB EDR and FPGA-direct
network• Some nodes are equippedwith both FPGAs and GPUs,
and other nodes are with
GPUs only

82022/05/25

PCクラスタOSSワークショップ（FPGA）

SINGLENODE(with FPGA)CPU

Network switch(100Gbps x2)

HCA

Network switch(100Gbps x2)

CPU

HCA

HCA

PCIe network (switch)

PCIe network (switch)

G
P
U

G
P
U

G
P
U

HCA

G
P
U

Center for Computational Sciences, Univ. of Tsukuba

Two types of interconnection network
Inter-FPGA direct network(only for Albireo nodes)
FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

InfiniBand HDR100/200 network for parallel processing
communication and shared file system access from all nodes

IB HDR100/200 Network (100Gbps x4/node)
…

FPGA

FPGA

comp.
node

…

comp.
node

Deneb nodes
64 of FPGAs on Albireo nodes (2FPGAS/node)
are connected by 8x8 2D torus network
without switch

92022/05/25

PCクラスタOSSワークショップ（FPGA）

comp.
node

comp.
node

Albireo nodes

comp.
node

…

comp.
node
…

Inter-FPGA torus network

For all computation nodes (Albireo and Deneb) are connected by full-bisection
Fat Tree network with 4 channels of InfiniBand HDR100 (combined to HDR200
switch) for parallel processing communication such as MPI, and also used toaccess to Lustre shared file system.

Center for Computational Sciences, Univ. of Tsukuba

G
P
U

G F
P P
U G
A

F G
P P
G U
A

G
P
U
IB HDR100 x4⇨ HDR200 x2

CPU

CPU
100Gbps x4
FPGA optical
network x2
IB HDR200
switch (forfull-bisection
Fat-Tree)

1.2Tbps/nodeAlbireo node10

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukubaこれまでの研究
nFPGA-network: CIRCUS (Communication Integrated Reconfigurable CompUting System)nnnGPU-FPGA DMA: FPGAから起動しCPUの助けを借りないnn

PCIeのプロトコルを利用したDMAでGPUメモリとFPGAメモリ（global memory）間で高速データ転送

Programming:nnn

FPGAボード間を光リンク（最大4本、100Gbps）で結合するネットワークを構築、ルータ機能をOpenCLから利用可能とする
演算と通信を細粒度でパイプライン処理 ⇨ FPGAの特性を並列処理に最大限に活かす

MHOAT (Multi-Hetero OpenACC Translator)⇒ OpenACCによる単一コード記述でGPUとFPGAのオフローディングを簡潔に記述
Intel oneAPI⇒ taskベースでGPUとFPGAのカーネル起動を簡潔に記述、DPC++を基本とするがOpenCLやCUDAも吸収可能

Appllication: 宇宙物理学コードARGOTn11

2種類の異なる動作と特性を持つ計算部分をGPUとFPGAに分散

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

CIRCUSnIntel FPGA SDK for OpenCLnn

FPGA hardware の起動を OpenCL APIで可能に問題：光リンクはあるがそのルーティングを全部OpenCLで記述するのか？
高速・低レイテンシの通信には Verilog HDL などのHDLによる記述が必要
n MPI等は memory-to-memory 通信モデルでFPGAにはなじまない
n FPGAの持つパイプライン演算特性を通信に延長
nn→CIRCUS: Communication Integrated Reconfigurable CompUting Systemn

演算と通信をシームレスに結合し、細粒度処理を実現

Comm.
receiver code on FPGA2
Backendreceiver(__global float* restrict x, int n) {sender(__global float* restrict x, int n)for (int i = 0; i < n; i++) {{float v = read_channel_intel(simple_in);for (int i = 0; i < n; i++) {x[i] = v;
float v = x[i];}write_channel_intel(simple_out, v);}}}

sender code on FPGA1

12* N. Fujita, et al., “Performance Evaluation of Pipelined Communication Combined with Computation in OpenCL Programming on FPGA”, AsHES2020.
2022/05/25 PCクラスタOSSワークショップ（FPGA）

max. throughput
90.2Gbps
min. latency
500ns

Throughput（1hop~7hops）

BetterBetterCIRCUS 通信性能（Cygnus）

latency+ /hop~250 ns

Latency（1hop~7hops）

Evaluated on up to 8 Bittware 520N FPGA boards in Cygnus supercomputer at CCS, University of Tsukuba
N. Fujita, et al., “Performance Evaluation of Pipelined Communication Combined with Computation in OpenCL Programming on FPGA”, AsHES2020.

13

2022/05/25

PCクラスタOSSワークショップ（FPGA）

演算・通信パイプライン融合による collective 通信n

CIRCUSの機能を使って隣接通信と加算をword単位でパイプライン処理

142022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

GPU-FPGA DMA

CPU

FPGA

15 2022/05/25

GPU

PCクラスタOSSワークショップ（FPGA）__kernel void fpga_dma(__global float *restrict fpga_mem,
const ulong gpu_memadr,
const uint id_and_len){cldesc_t desc;// DMA transfer GPU -> FPGA
desc.src = gpu_memadr;
GPU-to-FPGA DMA kickdesc.dst = (ulong)(&fpga_mem[0]);desc.id_and_len = id_and_len;
write_channel_intel(fpga_dma, desc);ulong status = read_channel_intel(dma_stat);}Center for Computational Sciences, Univ. of Tsukuba

FPGA-GPU DMA（Intel A10 + NVIDIA V100)
8

via CPU (FPGA ← GPU)
via CPU (FPGA → GPU)

7

Minimum latency (µsec)

Bandwidth [GB/s]

Better6

direction

54

via CPU

FPGA-GPU DMAGPU→FPGA

17

1.44

FPGA→GPU20

0.60[Reference]

3

•

2
1
0
1

16

256

4096

65536

Ryohei Kobayashi, Norihisa Fujita,
Yoshiki Yamaguchi, Ayumi Nakamichi,
Taisuke Boku, "GPU-FPGA
Heterogeneous Computing with
OpenCL-enabled Direct Memory Access",
Proc. of Int. Workshop on Acceleratorsand Hybrid Exascale Systems(AsHES2019) in IPDPS2019 (to be
published), May 20th, 2019.

1048576

Size [Bytes]
PCクラスタOSSワークショップ（FPGA）

16 2022/05/25

Center for Computational Sciences, Univ. of Tsukuba

OpenACCによる統一的記述を実現するメタコンパイラ
n

OpenACC はNVIDIA GPUで使えるだけでなくFPGAでも利用可能（研究ベース）
筑波大学CCS、ORNL、理研R-CCSでメタコンパイラを共同開発
n GPU – PGI OpenACC コンパイラ
n FPGA– OpenARC for FPGA on OpenACC: OpenACCコードをIntel FPGA向けのOpenCLに変換nホストコード上のランタイムシステムとの細かいコンフリクトを解消し、両コンパイラを機能結合するトランスレータ⇨ MHOAT (Multi-Hybrid OpenACC Translater)
n 問題
n

n

n
n
n

17

GPUとFPGAの並列処理特性の違い

GPU – SIMD的な水平型データ並列処理
FPGA – クロックレベルでのパイプライン処理（＋マルチインスタンス）
メモリモデルの違い

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

Compilation flow of MHOATBackend (existing compilers)

single program

Object Files

OpenACC

openacc.c
int main() {~~~~~~~~~~#pragma acc ~~~=gpu~~~~~~~~~~~~~~~~~~~~#pragma acc ~~~=fpga~~~~~~~~~~~~~~~~~~~~PGI Compiler

gpu.c
M
H
O
A
T

fpga
.cpp

Executable
pgc++

OpenCL hostfpga.c~~~~~~~~~~
return 0;}

gpu.o(pgc++ -c)

pgc++ cfpga.o

Open
ARC

a.out

transfer
FPGA
binary& runtimeIntel FPGA SDK for OpenCL

MHOAT is under development
targeting the first app. of ARGOT
182022/05/25

PCクラスタOSSワークショップ（FPGA）

OpenCLkernelfpga.claoc -c

fpga
.aoco

aoc

fpga
.aocx

Center for Computational Sciences, Univ. of Tsukuba

MHOATの現状
簡単なコードは実行可能
n 実アプリケーション（ARGOT）のコンパイルではOpenARCコンパイラの特異性の解消が難しく細かいソースコード書き換えが必要
n ARGOTコードのコンパイルはできるが性能、演算結果に問題がある
n⇨ 開発を継続中

19

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

もう一つのアプローチ：Intel oneAPIn oneAPI is a cross-architecture

programming framework− simplify the development acrossdifferent architecturesn Data Pallalell C++ (DPC++)−DPC++ = C++ + SYCL +extensions−enables unified description acrossdifferent architectures
こちらについては2022年4月のPCクラスタコンソーシアムOSSワークショップで紹介済み
20

2022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

アプリケーション例n

ARGOT (Accelerated Radiative transfer on Grids using Oct-Tree)
Simulator for early stage universe where the first stars and galaxies were bornn Radiative transfer code developed in Center for Computational Sciences (CCS),
University of Tsukuba
n CPU (OpenMP) and GPU (CUDA) implementations are available
n Inter-node parallelisms is also supported using MPInnART (Authentic Radiation Transfer) method
It solves radiative transfer from light source spreading out in the spacen Dominant computation part (90%~) of the ARGOT programnn

We accelerate the ART method on an FPGA using Intel FPGA SDK for
OpenCL as an HLS environment (with oneAPI)
21

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

Cosmic Radiative Transfer Simulation
ARGOTコードの2つの計算要素：ARGOT法とART法*
ARGOT• ARGOT method: Point Source processing• ART method (Authentic Radiation Transfer): Diffused Photon processing

Point Source

22 2022/05/25 PCクラスタOSSワークショップ（FPGA）

Diffuse PhotonCosmic Radiative Transfer Simulation
ARGOTコードの2つの計算要素：ARGOT法とART法*
ARGOT• ARGOT method: Point Source processing• ART method (Authentic Radiation Transfer): Diffused Photon processing

ARGOT (Accelerated Radiative transfer on Grids using Oct-Tree) code
ARGOT scheme

ART scheme

for radiative transfer (RT)from point source

for RT from mattersspatially spreading out

GPU

GPU

GPU

CHARM
Point Source
GPU acceleration

Diffuse Photon

FPGA

23 2022/05/25 PCクラスタOSSワークショップ（FPGA）

FPGA

FPGA

FPGA acceleration

ARGOT法(CUDAによる実装)

n ARGOT法は点光源の輻射輸送を計算

target mesh gridn 八分木を用いて3次元空間に分散する点

光源を表す−遠距離にある点光源の集合を単一の光源とみなせる
計算量: O(N^2) → O(NlogN)n 重力計算における Tree-Code に似た手法
−GPU実装に適している

24 2022/05/25

24
PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

ART法(OpenCLによる実装)
n ART法は空間に広がる光源からの輻射輸送を計算−問題空間の端からレイが並行に直進ART法の実行時間は全体の90%を占める

n ART法はGPUに不向きな計算−レイの進行方向によってメッシュデータのメモリアクセスパ
ターンが変化
ランダムアクセスに近いアクセスパターン−Atomic演算が必要−計算中に含まれる演算数が不十分
n FPGAを用いて高速化−FPGAのon-chipメモリを活用−一定サイズのベクトル処理が独立に大量に必要
GPUの大規模SIMDが有効活用できない
詳細：Norihisa Fujita, et.al., “Accelerating Space Radiative Transfer on FPGA using OpenCL”, HEART2018

25

2022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

GPU-only vs GPU-FPGA 協調計算（ARGOT全体）

R. Kobayashi, et. al., “Accelerating Radiative Transfer Simulation with GPU-FPGA Cooperative Computation”, ASAP2020, Jul. 2020.

26

2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

OpenCL kernels.
ALM
695,204(74.5%)

Reg.
1,324,993(35.5%)

M20K
5,053(43.1%)

MLAB
1,916(33.3%)

Fmax[MHz]
206.67

ART法の並列FPGA性能

Results of performance evaluation (table). Eachin milliseconds.
1 Node2 Nodes
4 Nodes

GPU
174.9
268.3
355.3

FPGA
30.7
31.9
33.4

Breakdown of the GPU evaluation result. Eachin milliseconds.

e
e
e

GPU
Comp.
H2D
D2H
MPI
Others
Total
Figure 11: Parallel efficiency of the implementations at w
139.4
33.4
0.0 Weak0.0scalingによる4ノードまでのART法処理時間（左）と並列化効率（右）
2.1
174.9
scaling.
139.9 10:35.0Results
39.3 of performance
51.9
2.2 evaluation
268.3
Figure(chart).
*Norihisa 33.3
Fujita, et.al.,"OpenCL-enabled
Parallel
140.2
40.8
138.1
2.8 Raytracing
355.3 for Astrophysical Application on Multiple FPGAs with Optical Links", 2020 IEEE/ACM
International Workshop on Heterogeneous High-performance Reconfigurable Computing (H2RC) in conjunction with SC20, Nov. 2020.

changed when increasing the number of nodes. Accord
it is operated for practical use.1 We expect the performances
2022/05/25 PCクラスタOSSワークショップ（FPGA）
Center for Computational
Sciences, Univ.
Tsukuba
to Figure 10, the computation
time is constantifofweincre
the27
CygnusPPX FPGAresults.
to be thesame because
10ofand
Table
IVFPGA
show andthe evaluation
These
theyinclude
are compatible
terms offorthethesystem
hardwarethe
andnumber of nodes.
notthe timeinrequired
data copy
As discussed in the previous subsection, it is diffi
FPGA-relatedheinitial valuesoftware.
and retrieve the result. Table V
for us to evaluate the communication time in the FP

oneAPI (single node) 性能評価
n 問題サイズ: 32!

n Others−ベクトル並列化が可能
GPUによる高速化が効果的
n ART

Better−ART法のGPU実装はCPUと同程度の性能−FPGAにより高速化
n GPU:CUDA+FPGA:OpenCL

実行時間は1.5%だけ増加−GPUのみの場合と比べて10xの性能向上
n oneAPI vs CUDA+OpenCL−oneAPI版の実行時間が1.5%増加→oneAPI処理のオーバヘッドは無視できる

28

2022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

ARGOT法およびART法の非同期実行
n 本来ARGOT法とART法は2種類の物理現象を独立に計算し、最後に結果を足し合わせている

0.7
Othersn ART法の実行時間は ARGOT法の実行時間に隠蔽される−全体として1.38xの高速化

Lower is better−OpenMPを利用して，各Queueにactionを非同期に投入−Queueに投入されたactionは非同期実行される

Execution time [s]n 各Queueは非同期に実行可能

0.6

ARGOT

0.5

ART

0.4
0.3
0.2
0.1
0
ARGOT(GPU)ART(FPGA) / CUDA + OpenCL

ARGOT(GPU)ART(FPGA) / CUDA +OpenCL over oneAPI

ARGOT(GPU)ART(FPGA) / CUDA +OpenCL over oneAPI async*柏野隆太他, “oneAPIを用いたGPU・FPGA混載ノードにおける宇宙物理シミュレーションコードARGOTの実装”, 第183回HPC研究会

292022/05/25

PCクラスタOSSワークショップ（FPGA）
Center for Computational Sciences, Univ. of Tsukuba

まとめと今後の展望
n

FPGA単体でのHPC利用はなかなか難しいnnn

絶対性能（FLOPS）
メモリテクノロジ

GPUとFPGAを相補的に利用することでmulti-physics simulationに柔軟に対応
nnn

バルク並列処理
複雑な処理、条件分岐
ノード間通信n

筑波大学CCSではCHARM (Cooperative Heterogeneous Acceleration with Reconfigurable Muti-devices)コンセプトの下、GPUとFPGAを適材適所で用いることで全方位的な新しいHPC solutionを開拓している
Big Data/AI処理のベースは現状ではDeep LearningでGPUがそれに特化されつつあるが、柔軟な部分処理にはFPGAが適用できる（例：sort engine、staging処理等）
HPC/AIのエンドユーザ向けのプログラミングが最大の問題

n

今後の展望

n

n

n
n
n
n

30

CygnusにおけるoneAPIの本格利用
ORNL-OpenARCとは独立に理研R-CCSと筑波大の共同研究で Omni OpenACC-OpenCL translater (for FPGA)を開発中
MHOATの実装とARGOTのような実コードの高速化
富岳/ESSPERにARGOTを移植中
2022/05/25

PCクラスタOSSワークショップ（FPGA）

Center for Computational Sciences, Univ. of Tsukuba

研究協力
n

筑波大学・計算科学研究センター/筑波大学・システム情報系

31

小林涼平

藤田典久

梅村雅之

吉川耕司

2022/05/25

PCクラスタOSSワークショップ（FPGA）

山口佳樹

Center for Computational Sciences, Univ. of Tsukuba