Multi-Hybrid Accelerated Supercomputer:
The new system of CCS
Taisuke Boku
HPC System Division, Center for Computational Sciences
University of Tsukuba
A part of slides are based on research collaboration with
N. Fujita, Y. Yamaguchi, R. Kobayashi, Y. Ohbata, M. Umemura and K. Yoshikawa

1

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

Outline
Supercomputer history of CCSn Next generation’s accelerated computing systemn FPGA for HPC as large scale parallel systemn AiS (Accelerator in Switch) conceptn PACS-X Project and PPXn Ongoing research on PPX
n Brief summary of new machinen Summaryn2

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Our History and Current Machines

3

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

History of PAX (PACS) series at U. Tsukuba
n
n
n

1977: research started by T. Hoshino and T. Kawai
1978: PACS-9 (with 9 nodes) completed
1996: CP-PACS, the first vendor-made supercomputer at CCS, ranked as #1 in TOP500
1978

1st gen: PACS-9

Year
1978
1980
1983
1984
1989
1996
2006
2012~13
2014

1996

1980

2nd gen. PACS-32

1989
5th gen, QCDPAX

Name
PACS-9PACS-32
PAX-128
PAX-32JQCDPAX
CP-PACS
PACS-CS
HA-PACSCOMA (PACS-IX)

6th gen: CP-PACS
Ranked #1 in TOP500

Performance
7 KFLOPS
500 KFLOPS
4 MFLOPS
3 MFLOPS
14 GFLOPS
614 GFLOPS
14.3 TFLOPS
1.166 PFLOPS
1.001 PFLOPS

2006

2012~2013

7th gen: PACS-CS

8th gen: GPU cluster HA-PACSn

co-design by computer scientists
and computational scientists
toward “practically high speed
computer”n

Application-driven developmentn

Sustainable development
experience

CCS  Symposium  2018,  Tsukuba

4

2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Two streams of supercomputers in CCS (~ FY2017)n

Peak performance for advanced users: HA-PACS (PACS-VIII)
GPU cluster + originally developed inter-GPU communication system TCA (Tightly Coupled
Accelerator)n Accelerated computing with PFLOPS class performance for code development and product runn Professional users rather than novice usersnn

Easy programming for general users: COMA (PACS-IX) (➙ Oakforest-PACS)
Many core architecture for general programming (OpenMP + MPI) toward next generation’s
many-core (Intel Knights Corner)n Leading system toward Oakforest-PACS (by JCAHPC: joint organization with U. Tsukuba and U.
Tokyo) which is equipped with Intel Knights Landingn Wide area and general users with ordinary programmingnn

HA-PACS retired on March 2018 and COMA will retire on March 20195

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Next Generation’s
Accelerated Computing

6

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

Accelerators in HPC
n

Traditionally...
Cell Broadband Engine, ClearSpeed, GRAPE. ...
n then GPU (most popular)nn

Is GPU perfect ?
good for many applications (replacing vector machines)n depending on very wide and regular computationnlarge scale SIMD (STMD) mechanism in a chipn high bandwidth memory (GDR5, HBM) and local memorynn

bad fornot enough parallelismn not regular computation (warp spliting)n frequent inter-node communication (kernel switch, go back to CPU)n

7

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

FPGA in HPC
nGoodness of recent FPGA for HPC
True codesigning with applications (essential)n Programmability improvement: OpenCL, other high level languagesn High performance interconnect: 40Gb~100Gbn Precision control is possiblen Relatively low powernn

Problems
Programmability: OpenCL is not enough, not efficientn Low standard FLOPS: still cannot catch up to GPU-> “never try what GPU works well on”n Memory bandwidth: 2-gen older than high end CPU/GPU-> be improved by HBM (Stratix10)n

8CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Simple pros/cons

performance(FLOPS)

external
communication (sec, B/s)

programming
cost

CPU

△

○

◎

GPU

◎

△

○

FPGA

○

◎

×➝△?

How to compensate with each othertoward large degree of strong scaling ?

9

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Multi-Hybrid Accelerated ComputingnCombining goodness of different type of accelerators: GPU + FPGA
GPU is still an essential accelerator for simple and large degree of parallelism to provide ~10
TFLOPS peak performancen FPGA is a new type of accelerator for application-specific hardware with programmability andspeeded up based on pipelining of calculationn FPGA is good for external communication between them with advanced high speed
interconnection up to 100Gbps x4chan.
nnMulti (two type of devices) – Hybrid (acceleration) Supercomputer
PC cluster with multiple acceleration devices: GPU + FPGA (+CPU)n Mixing their good feature toward high performance computational sciences especially for strong
scaling covering various size of computation on complicated calculation such as multi-physicalsimulationn Using FPGA’s external communication link for low latency and high bandwidth communicationn

10

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

AiS: conceptual model of Accelerator in Swtich
invoke GPU/FPGA kernsls
data transfer via PCIe
(invoked from FPGA)
GPU
CPU

PCIe

•
•

FPGA can work both forcomputation and communicationin unified manner
GPU/CPU can request
application-specificcommunication to FPGA

11

FPGA
comp.
comm.

CPU

PCIe
collective or specialized
communication

GPU
FPGA
comp.
comm.

>  QSFP+  interconnect

Ethernet Switch

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

PACS-X Project
toward next generation‘s
hybrid supercomputing

12

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

PACS-X (ten) Project at CCS, U. TsukubanHistory of supercomputers: PACS (PAX)a series of co-design base parallel system development both on system andapplication at U. Tsukuba (1978~)n recent systems focus on acceleratorsn

PACS-VIII: HA-PACS (GPU cluster, Fermi+Kepler, PEACH2, 1.1PFLOPS)n PACS-IX: COMA (MIC cluster, KNC, 1PFLOPS)nn

Next generation of TCA implementation
new generation of GPU and FPGA with high speed interconnectionn more tightly co-designing with applicationsn system deployment starts from 2017n

PPX: Pre-PACS-X
13

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

PPX: testbed under AiS concept (x 13 nodes)
1.6TB  NVMe

Xeon Broadwell
QPI

HCA:
Mellanox IB/EDR

100G IB/EDR
GPU: coarse-grain
offloading
NVIDIA P100 x 2

Xeon Broadwell

40Gb Ether x 2-> upgrade to 100G x 2
14

FPGA:
fine-grain
partial offloading+
high-speed interconnect
Intel Arria10 or Xilinx Kintex

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

PPX mini-cluster system
InfiniBand/EDR  Switch
InfiniBand/EDR (100Gbps)
computenode
CPU: BDW x 2
GPU: P100 x 2
FPGA x1

13 nodes....

Ethernet (40Gbps x 2)

100G  Ethernet  Switch
Ethernet will be upgraded to 100Gb x4 on next gen.

15

GbE switch

login node

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Current Research on PPX
toward PACS-X Real System16

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

OpenCL-enabled high speed networkn

OpenCL environment is availableex) Intel FPGA SDK for OpenCLn basic computation can be written in OpenCL without Verilog HDLnnBut, current FPGA board is not ready for OpenCL on interconnect accessnn

BSP (Board Supporting Package) is not complete for interconnect➝ we developed for OpenCL accessOur goalenabling OpenCL description by users including inter-FPGA communicationn providing basic set of HPC applications such as collective communication, basic linear libraryn providing 40G~100G Ethernet access with external switches for large scale systemsn

17

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Communication paths
IB (InfiniBand) Switch
IB EDR(100Gbps)

NODE

PCIe Gen.3x16

CPU0

via-IBIB HCA

CPU0

IB HCA
NODE

QPI

CPU1

FPGA

CPU1

・・・・

FPGA

PCIe Gen.3x8(56Gbps)
QSFP+(40Gbps)via-Ethernet
Ethernet Switch

18

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Communication latency
node-to-node communication
2経路間における通信レイテンシの比較
latency (1 byte data)(通信データサイズ: 1 byte)
29.03
FPGA<->CPU1

better

レイテンシ[us]Latency [usec]

25

27.70
FPGA -> CPU
CPU -> FPGAの合計

15
10

0

1.33
IB経由

via IB19

2018/10/16n

CPU-FPGA interface by current
BSP is not good

CPU1->CPU1

20

5

n

~1µs latency on Ethernet
CPU-FPGA comm. occupies
comm. latency via IB

0.99
Ethernet経由

via Ethernet

1.2

Latency [usec]
レイテンシ[us]

30

n

break
down of Ethernet latency
Ethernet経由のレイテンシの内訳

10.8

Ethernet IP
Controller

0.6

Ethernet IP

0.4

Ethernet
Switch0.2
0via Ethernet
Ethernet経由

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

Communication bandwidth• 40Gbps Ethernet achieves 4.97GB/s- 99.8 % of theoretical peak (w/o error handling)- small N1/2 by short latency

To be upgraded with 100Gbps x4chan. on
next generation Stratix10• via-IB achieves 2.32GB/s- non-pipelined- no special feature (such as GPUDirect) on FPGA-HCA
node-to-node communication
2経路間における通信バンド幅の比較
bandwidth

Bandwidth [GB/s]
バンド幅[GB/s]

better

6
5
4

IB経由

3

Ethernet経由

2
1
0

16
16

256
256

4096
4K

65536
1048576
64K
1M
データサイズ[byte]

16777216
16M

268435456
256M

Data size [Byte]

20

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Channel over Ethernet (CoE): new implementationn

ChannelExchange data between kernels in an FPGAn One of Intel FPGA’s extensions to the OpenCL languagenn

Channel over Ethernet
enables us transfer data in channel over the networkn Using same APIs as Channel from the user coden Ethernet as a protocol on the networkn

sender code on FPGA1__kernel  void  sender(__global  float*  restrict  x,  int n)  {receiver code on FPGA2
for  (int i =  0;  i <  n;  i++)  {float  v  =  x[i];Eth. __kernel  void  receiver(__global  float*  restrict  x,  int n)  {write_channel_intel(simple_out,  v);for  (int i =  0;  i <  n;  i++)  {}float  v  =  read_channel_intel(simple_in);}x[i]  =  v;}
CCS  Symposium  2018,  Tsukuba
21}
Center for Computational Sciences, Univ. of Tsukuba
2018/10/16

Enabling Networking from OpenCL
Board Support Package (BSP) is
FPGA board (A10PL4)
a hardware component
Ethernet IP
describing a board specification FPGA
Controller
DDR4n Adding network controller into
Controller
OpenCL
DDR4
BSP
kernelInterconnectn

40Gb Ethernet IPn Bridge logic between OpenCL
kernels and the IPnn

OpenCL kernels can use it
through I/O Channels.
22

Controller
PCIe
Controller

Ethernet
IP Core

QSFP+
DDR4
mem
DDR4
mem

BSP

DriverHost PC
Host Application

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Implementation[I+]

expanderpacketbuffer[I-]

expanderpacketbuffer[J+]

expanderpacketbuffer[J-]

expanderpacketbuffer[stage1]

scalar[stage2]

scalar[stage3]

scalar[stage4]

scalarHimenoPEn

Packet handler logics are written in
OpenCL
Buffering and binary merge tree for sendingn Decoding packets for receivingn Tree algorithm for allreducenn

We decide to use this design to reduce
modifying cost of BSP
n

Interface to the BSP is fixednn

can be used for a variety of applications

But resource usage may be increased

Allreduce

packet (280bit)
packed data (136bit)
input data (32bit)

binary tree merge network

p0

p1

p2

p3

=

=

=

=

=

=

+

+

+

+

+

+
External Link

23

CCS  Symposium  2018,  Tsukuba
2018/10/16

Internal Link
p0

p1

p2

p3

Center for Computational Sciences, Univ. of Tsukuba

Preliminary ResultnHimeno Benchmark [1]n

3D-Poisson Equation Solver, 19-point stencil computationn

Performance is limited by memory bandwidth

Pipeline implementation using OpenCL:
direct data transfer over the network to/from the pipelinen Size XS (33x33x65) on 1, 2 and 4 FPGAs (Strong Scaling)nn

XS 1 node @210MHz

XS 2 nodes @225MHz

XS 4 nodes @208MHz

5019 MFLOPS

9762 MFLOPS

15225 MFLOPS

x1.00

x1.95

x3.03

Scaling seems to be struggled by the problem sizen

XS has relatively larger halo size than other sizes.
CCS  Symposium  2018,  Tsukuba[1]:
24 http://accc.riken.jp/en/supercom/himenobmt/
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

GPU-FPGA direct communication
On PCIe (gen3 x16 lanes), FPGA and GPU can directly
communication with DMA controller on FPGA toenhance the performance both on latency andbandwidthn Ordinary way: GPU->CPU, CPU->FPGA (and reversedway)n FPGA invokes DMA data transfer from/to GPU actively
during computation without support by CPU(initial parameters should be preset by CPU)n

25

GlobalMemoryDMA controller

A
Descriptor
controller* Hanawa, T et al., Interconnection Network for Tightly Coupled Accelerators,
2013 IEEE 21st Annual Symposium on High-Performance Interconnects, pp 79-82

Block
RAM

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

void*  d_m;  cudaMalloc(&d_m,  numdata*sizeof(uint32_t));  PCIe address mapping
unsigned long long paddr;
tcaCreateHandleGPU(&paddr,  d_m,  numdata*sizeof(uint32_t));if (m  ==  MODE_READ)  {  for (int n  =  1;  n  <=  maxnumdata;  n  *=  2)  {if (n  ==  maxnumdata)  n  =  numdata;
uint64_t count_sum =  0;
for (int i =  0;  i <  n;  ++i)  h_m[i]  =  n  +  i;
cudaMemcpy(d_m,  h_m,  n*sizeof(uint32_t),  cudaMemcpyHostToDevice);for (int t  =  0;  t  <  numtry;  ++t)  {descriptor transfer
int id  =  random_id();rd_ast_rx(ctrl,  paddr,  ONCHIP_MEMORY_OFFSET,  n*sizeof(uint32_t),  id);26

uint32_t status  =  rd_ast_tx(ctrl);
watch termination signal from CPU
if (status  ==  unsigned(0x0100)  +  id)  {count_sum +=  count_result_rd(ctrl);}  else {std::cerr <<  "incomplete  status  =  " <<  std::hex  <<  status  <<  std::endl;
exit(1);
Code for FPGA ← GPU}
data transfer
CCS  Symposium  2018,  Tsukuba}
2018/10/16

GPU-FPGA data transfer latencyn

Data size: 4 Bytesnn

FPGA ← GPUnn

Minimum size of data to be transferred
5.5x faster than ordinary method

MethodMin. latency

Via CPU

8.04 µsec

DMA1.45 µsec

MethodMin. latency

Via CPU

8.18 µsec

DMA

0.43 µsec

FPGA → GPUn

27

19x faster than ordinary method

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

GPU-FPGA data transfer bandwidth
Data size: 4 ~ 1M Bytesn

Current bandwidth is bound by PCIe for A10PL4 (Arria10) of PCIe gen.3 x8lanes
Via CPU (FPGA⇨GPU)
DMA (FPGA⇨GPU)

Bandwidth [GB/s]

Via CPU (FPGA⇦GPU)
DMA (FPGA ⇦ GPU)

Higher is BetternData size [Byte]
28

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

ARGOTn

ARGOT (Accelerated Radiative transfer on grids using Oct-Tree)
Radiative transfer code developed in Center for Computational Sciences (CCS),
University of Tsukuban CPU (OpenMP) and GPU (CUDA) implementations are availablen Inter-node parallelisms is also supported using MPInnART (Authentic Radiation Transfer) method
It solves radiative transfer from light source spreading out in the spacen Dominant computation part (90%~) of the ARGOT programnn

In this research, we accelerate the ART method on an FPGA using IntelFPGA SDK for OpenCL as an HLS environment

29

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

ART MethodnART method is based on ray-tracing method
n
n

n

n

ART method computes radiative intensity on each mesh as
shows as formula (1)
n
n

n

n

3D target space split into 3D meshes
Rays come from boundaries and move in straight in parallel witheach other
Directions (angles) are given by HEALPix algorithm

Bottleneck of this kernel is the exponential function (expf)
There is one expf call per frequency (ν). Number of frequency isfrom 1 to 6 at maximum, depending on the target problem
All computation uses single precision computations

Memory access pattern for mesh data is varies dependingon ray’s direction
n
n

30

Not suitable for SIMD style architecture
FPGAs can optimize it using custom memory access logics.

CCS  Symposium  2018,  Tsukuba
2018/10/16

Parallelizing ART method using channelsn

PE (Processing Element)

BE

BE

BE

PE

PE

BE

BE

PE

PE

BE

Computation kernel of ART method
3n Each PE has 8 size memory of meshnn

BE (Boundary Element)
BE handles ray I/O including boundary processingn R/W access of boundary, or R/W access of ray
buffernnChannels are used for communication
Transferring ray data to/from neighbor PE or BEn Two channels are used for each connection
because a channel is one-sidedn

31

96bit x2
(read,write)
Channel

y
x

BE

BE

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Ray Data

14C: 14 CPU cores=1 socket
28C: 28 CPU cores=2 sockets

Performance (FPGA vs. CPU)

Better1282.8n
n

1165.21111.01133.5

4.6 times faster
112.4 (14C)

183.4 (28C) 227.2 (28C)

165.0 (28C)

FPGA is faster than CPU on all problem sizes
FPGA is 4.6 times faster on 643 and 6.9 times faster on 1283
n
n

32

Decreasing on CPU from 643 to 1283 is caused by cache missSize of mesh data cannot fit into the cache
CCS  Symposium  2018,  Tsukuba
2018/10/16

Cygnus
New Supercomputer at CCS

33

2018/10/16

CCS  Symposium  2018,  Tsukuba
Center for Computational Sciences, Univ. of Tsukuba

Objective and time line
Since COMA system will retire in the end of March 2019, we need a new machine as afollow-up system from April 2019n We decided to deploy a new system with CCS budget with partial support from MEXTn The new system is designed based on AiS concept to realize Multi-Hybrid Accelerated
Computing with GPU and FPGAn Time Line
n

n
n
n
n
n
n
n
n

34

RFI: Feb. 21st 2018
RFC: Jun. 8th 2018
RFP: Aug. 9th 2018
Proposal submission: Sep. 25th 2018 ⇨ Final system specification fixing (not yet today!!)
Bid open: Oct. 25th 2018
Specification data shown today
Delivery: Mar. 15th 2019is subject to be changed
Test operation: Apr. 1st 2019
Real program starts: May 1st 2019CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

System name and constructionn

System name: “Cygnus”nn

based on “Cygnus A” galaxy with two accelerated radiation streams

Consisting of two parts
GPU-only part: equipped with CPU and GPU to be used as ordinary GPU cluster ⇨ Deneb?
n GPU+FPGA part: equipped with GPU and FPGA connected via PCIe and CPU to be used as AiS
system, as well as ordinary GPU cluster (not using FPGA) ⇨ Albireo?
n CPU and GPU of either part are identical as well as PCIe switch configurationn

Two parts are combined with InfiniBand high performance interconnectionn When FPGA-ready job is not running, GPU+FPGA part is also available as GPUonly part for high utilization ratio of the systemn Each computation node of either part is with “fat” configuration:very high performance both on computation and communicationn

35

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Single node configuration (with FPGA): “fat node”• 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

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)

36

Network switch(100Gbps x2)

G
P
U

FPGA

Inter-FPGA
direct network(100Gbps x4)

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Single node configuration (without FPGA): “fat node”• 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

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)

37

Network switch(100Gbps x2)

G
P
U

FPGA

Inter-FPGA
direct network(100Gbps x4)

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Approximate specification of Cygnus

38

Item

Specification

Peak performance~3.0 PFLOPS DP(GPU: 2.3 PFLOPS, CPU: 0.12 PFLOPS, FPGA: 0.6 PFLOPS)⇨ enhanced by mixed precision and variable precision on FPGA

# of nodes~80 (~32 GPU+FPGA nodes, ~48 GPU-only nodes)

CPU / nodeIntel Xeon Gold x2 sockets

GPU / node

NVIDIA V100 x4 (PCIe)FPGA / nodeIntel Stratix10 x2 (each with 100Gbps x4 links)

Global File System

Lustre, RAID6, ~3.0 PB

Interconnection Network

Mellanox InfiniBand EDR (HDR100) x4

Total Network B/W~4 TB/s

Programming Language

CPU: C, C++, Fortran, OpenMP
GPU: OpenACC, CUDA
FPGA: OpenCL, Verilog HDL

CCS  Symposium  2018,  Tsukuba
2018/10/16

Two types of inter-node network

FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

FPGA

IB EDR Network (100Gbps x4/node)

comp.
node

comp.
node

…

comp.
node

Ordinary inter-node communication channel for
CPU and GPU, but they can request it to FPGA

(only for GPU+FPGA node)
39

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

High Level Programming Paradigmn

XcalableACC
n
n

n
n

n

under development in collaboration between CCS-Tsukuba and RIKEN-AICS
PGAS language XcalableMP is enabled to imply OpenACC for sophisticated coding of distributed
memory parallelization with accelerator
inter-node communication among FPGA can be implemented by FPGA-Ethernet direct link
Data movement between GPU and FPGAOpenACC for FPGA
n
n
n

40

(plan) research collaboration with ORNL FTG
OpenACC -> OpenCL -> FPGA compilation by OpenARC project is under development
final goal: XcalableACC with OpenARC compiler and FPGA-Ethernet link

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba

Summary
CCS will introduce a new supercomputer based on AiS concept as the next
generation‘s Multi-Hybrid Accelerated Supercomputern New machine Cygnus will be equipped with very high performance GPU and
FPGA (partially) to make “strong scaling ready” accelerated system forapplications where GPU-only solutions is weak, as well as all kind of GPUready applicationsn FPGA for HPC is a new concept toward next generation’s flexible and low
power solution beyond GPU computingn Cygnus system will be delivered on next March and official program start on
May 2019n

41

CCS  Symposium  2018,  Tsukuba
2018/10/16

Center for Computational Sciences, Univ. of Tsukuba