with Proprietary Interconnect and Its Programming and Applications - - PowerPoint PPT Presentation
Tightly Coupled Accelerators with Proprietary Interconnect and Its Programming and Applications Toshihiro Hanawa Information Technology Center, The University of Tokyo Taisuke Boku Center for Computational Sciences, University of Tsukuba
GPU Technology Conference 2015
1
Background HA-PACS / AC-Crest Project Introduction of HA-PACS / TCA
Organization of TCA PEACH2 Board designed for TCA
Evaluation of Basic
Collective Communications
Implementation Examples Performance Evaluation
Application Examples
QUDA (QCD) FFTE (FFT)
Introduction of XcalableACC
Concept Code Examples Evaluations
Summary
GPU Technology Conference 2015
2
Advantageous Features
Examples of HPC System:
GPU Clusters and MPPs in TOP500
2nd: Titan (NVIDIA K20X, 27 PF) 6th: Piz Daint (NVIDIA K20X, 7.8 PF) 10th: Cray CS-Storm (NVIDIA K40, 6.1
15th: TSUBAME2.5 (NVIDIA K20X, 5.6
48 systems use NVIDIA GPUs.
GPU Clusters in Green500 (Nov.
3rd: TSUBAME-KFC (NVIDIA K20X, 4.4
4th: Cray Storm1 (NVIDIA K40, 3.9 GF /
7th: HA-PACS/TCA (NVIDIA K20X, 3.5
8 systems of Top10 use NVIDIA GPUs.
GPU Technology Conference 2015
3
Ultra-low latency between GPUs is important for next generation’s HPC
GPU Technology Conference 2015
4
8th generation of PAX/PACS series supercomputer in University of Tsukuba
FY2011-2013, operation until FY2016(?)
Target field: QCD, astrophysics, QM/MM (quantum mechanics / molecular mechanics, bioscience)
for development of GPU-accelerated code for target fields, and performing product-run
Now in operation since Feb. 2012
for elementary research on direct communication technology for accelerated computing
Our original communication chip named “PEACH2” was installed in each node.
Now in operation since Nov. 2013
GPU Technology Conference 2015
5
Project “Research and Development on Unified Environment of
GPU Technology Conference 2015
6 Objectives
Higher productivity Highly optimized (offload, communication)
Eliminate extra memory copies to the host Reduce latency, improve strong scaling with small
Enable hardware support for complicated
GPU Technology Conference 2015
7
8
CPU PCIe Switch Node CPU Memory
PCIe
GPU GPU Memory
PCe
CPU PCIe Switch Node CPU Memory
PCIe
GPU GPU Memory
PCe
PCIe
PEACH2 PEACH2
Using PCIe as a communication link
Direct device P2P communication is
PEACH2:
GPU Technology Conference 2015
Data copy between CPU and GPU (1 and 3) have to perform
GPU Technology Conference 2015
9
PCIe SW
PCIe SW
Temporal data copy is eliminated Lower latency than the previous method Protocol conversion is still needed
10
PCIe SW
PCIe SW
GPU Technology Conference 2015
direct data copy using GDR much lower latency than InfiniBand
11
PCIe SW
PCIe SW
GPU Technology Conference 2015
PEACH2 can access all
NVIDIA Kepler architecture
Connect among 3 nodes
Similar to ordinary GPU cluster
80 PCIe lanes are required
GPU Technology Conference 2015
(Xeon E5 v2)
(Xeon E5 v2)
GPU GPU 2 GPU 3
IB HC A
PEA CH2 GPU 1
G2 x8 G2 x16 G2 x16 G3 x8 G2 x16 G2 x16 G2 x8 G2 x8 G2 x8
QDR 2port
12
Performance over QPI is
PEACH2 is available for GPU0,
Note that InfiniBand with GPU
Similar to ordinary GPU cluster
80 PCIe lanes are required
GPU Technology Conference 2015
(Xeon E5 v2)
(Xeon E5 v2)
GPU GPU 2 GPU 3
IB HC A
PEA CH2 GPU 1
G2 x8 G2 x16 G2 x16 G3 x8 G2 x16 G2 x16 G2 x8 G2 x8 G2 x8
13
Implement by FPGA with four
Altera Stratix IV GX Prototyping, flexible
Sufficient communication
PCI Express Gen2 x8 for each
Sophisticated DMA controller
Chaining DMA, Block-stride transfer
function
Latency reduction
Hardwired logic Low-overhead routing
Efficient address mapping in
Simple comparator for decision
GPU Technology Conference 2015
14
GPU Technology Conference 2015
15
Main board + sub board
Most part operates at 250 MHz (PCIe Gen2 logic runs at 250MHz) PCI Express x8 card edge
Power supply for various voltage
DDR3- SDRAM
FPGA (Altera Stratix IV 530GX)
PCIe x16 cable connecter PCIe x8 cable connecter
GPU Technology Conference 2015
16
GPU Technology Conference 2015
17
Base cluster (Feb. 2012) TCA (Nov. 2013) Node CRAY GreenBlade 8204 CRAY 3623G4-SM MotherBoard Intel Washington Pass SuperMicro X9DRG-QF CPU Intel Xeon E5-2670 x 2 socket (SandyBridge-EP, 2.6GHz 8 core) x2 Intel Xeon E5-2680 v2 x 2 socket (IvyBridge-EP, 2.8GHz 10 core) x2 Memory DDR3-1600 128 GB DDR3-1866 128 GB GPU NVIDIA M2090 x4 NVIDIA K20X x 4 # of Nodes (Racks) 268 (26) 64 (10) Interconnect Mellanox InfiniBand QDR x2 (Connect X-3) Mellanox InfiniBand QDR x2 + PEACH2 Peak Perf. 802 TFlops 364 TFlops Power 408 kW 99.3 kW
GPU Technology Conference 2015
18
(2.8 GHz x 8 flop/clock)
Total: 5.688 TFLOPS
8 GB/s
AVX
1.31 TFLOPSx4 =5.24 TFLOPS 22.4 GFLOPS x20 =448.0 GFLOPS
(16 GB, 14.9 GB/s)x8 =128 GB, 119.4 GB/s (6 GB, 250 GB/s)x4 =24 GB, 1 TB/s 4 Channels 1,866 MHz 59.7 GB/sec 4 Channels 1,866 MHz 59.7 GB/sec
Ivy Bridge Ivy Bridge
4 x NVIDIA K20X
Gen 2 x 16 Gen 2 x 16 Gen 2 x 16 Gen 2 x 16
PEACH2 board (Proprietary Interconnect for TCA)
Gen 2 x 8 Gen 2 x 8 Gen 2 x 8
Red: upgraded from base-cluster to TCA Legacy Devices
GPU Technology Conference 2015
19
GPU Technology Conference 2015
20
LINPACK: 277 Tflops (Efficiency 76%) 3.52GFLOPS/W #3 Green500 at Nov. 2013
Each group consists of 2
Orange: Ring Red: Cross link between 2
In TCA sub-cluster, 32 GPUs
limited to 2 GPUs under the
GPU Technology Conference 2015
21
DMA controller function
Chaining
Multiple DMA descriptors chained
DMA transactions are automatically
Block-stride support DMA Engine with 4ch
GPU Technology Conference 2015
22
Comm. type
Min. Latency Band width How working Comm. patterns
DMA Low
(< 2us)
High
DMA controller Any CPU
PIO Very low
( < 1us)
Low
CPU’s write
CPU-CPU
Source Destination Length Flags Next Source Destination Length Flags Next Source Destination Length Flags Next = NULL Head of Descriptors DMAC Control Reg.
Ping-pong performance between
Latency and bandwidth Written as application
Comparison
MVAPICH2-GDR 2.0b (with/without
A InfiniBand QDR link (40Gbps) is used,
which has the same performance as PEACH2.
Performance over QPI on TCA nodes
Special driver named “TCA p2p
GPU Technology Conference 2015
23
1 2 3 4 5 6 7 8 9 10 8 64 512 4096 32768
PIO: CPU to CPU: 0.8 us DMA:CPU to CPU: 1.8 us
GPU Technology Conference 2015
24
PIO: CPU to CPU: 0.8 us DMA:CPU to CPU: 1.8 us
200-300 nsec
BW converges to the same peak
GPU Technology Conference 2015
25 1 2 3 4 5 6 7 8 8 64 512 4096 32768
Latency (usec) Data Size (bytes)
DMA Direct DMA 1 hop DMA 2 hop DMA 3 hop
500 1000 1500 2000 2500 3000 3500 4000 8 128 2048 32768 524288 Bandwidth (MB/sec) Size (Bytes)
95% of theoretical peak
Converge to the same peak if hop count increases
better than MV2GDR under < 1MB
Over QPI: limited to 360MB/s
SB(SandyBridge): limited to 880MB/s due to PCIe sw perf.
GPU Technology Conference 2015
26
Max Payload Size = 256byte Theoretical peak (detailed):
4GB/sec × 256 / (256 + 24) = 3.66 GB/s
DMA (SB, GPU)
DMA (QPI, GPU)
GPU Technology Conference 2015
27
Allgather
All processes gather data of each
Communication bandwidth as well
GPU-GPU DMA
Allreduce
Conduct specified operation
Latency decides the performance. CPU-CPU PIO with host copy
Alltoall
All processes exchange specific
Communication bandwidth is
GPU-GPU DMA, all requests to
GPU Technology Conference 2015
28
29
Requires (log2p) steps
Ex. p=16 => 4 steps
Node mapping optimization
1.
2.
GPU Technology Conference 2015
30
Requires (log2p) steps
Ex. p=16 => 4 steps
Node mapping optimization
1.
2.
GPU Technology Conference 2015
31
Requires (log2p) steps
Ex. p=16 => 4 steps
Node mapping optimization
1.
2.
GPU Technology Conference 2015
32
Requires (log2p) steps
Ex. p=16 => 4 steps
Node mapping optimization
1.
2.
GPU Technology Conference 2015
33
Requires (log2p) steps
Ex. p=16 => 4 steps
Node mapping optimization
1.
2.
GPU Technology Conference 2015
34
Ring Neighbor Exchange Recursive Doubling Dissemination MPI
GPU Technology Conference 2015
Allreduce time for 8 Bytes scalar data
Di
TCA is more than twice
Low latency of TCA works effectively
35
Ring Neighbor Exchange Recursive Doubling Dissemination MPI
GPU Technology Conference 2015
QUDA: The open source Lattice QCD library
widely used as a LQCD library for NVIDIA GPUs
Optimized for NVIDIA GPUs All calculation run on GPUs
Solves a liner equation using CG method inter-node parallelism support supports multiple GPUs in a node source code is available at github
https://github.com/lattice/quda
36
GPU Technology Conference 2015
Write data to neighbor
latency is important
37
Halo data ＋ ＋ ＋ ＋
Allreduce Halo data exchange
GPU Technology Conference 2015
38
200 400 600 800 1000 1200 1400 MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA (2,1) (1,2) (4,1) (2,2) (1,4) (4,2) (2,4) (4,4) 2 Nodes 4 Nodes 8 Nodes 16 Nodes Time per iteration [us] Calc. Allreduce Comm.
(x,y) nodes
4 nodes configuration is the crossover point.
GPU Technology Conference 2015
39 (x,y) nodes
200 400 600 800 1000 1200 MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA MPI-P2P MPI-RMA TCA (2,1) (1,2) (4,1) (2,2) (1,4) (4,2) (2,4) (4,4) 2 Nodes 4 Nodes 8 Nodes 16 Nodes Time per iteration [us] Calc. Allreduce Comm.
GPU Technology Conference 2015
TCA: Tightly Coupled Accelerators
TCA enables direct communication among accelerators as an element technology becomes a basic technology for next gen’s accelerated computing in exa-scale era.
Bandwidth: max. 3.5 Gbyte/sec between CPUs (over 95% of theoretical peak), 2.8 Gbyte/sec between GPUs
CPUs), 2.0 us (DMA between GPUs)
GPU-GPU communication over the nodes can be utilized with 16 node sub-cluster.
Ping-pong program: PEACH2 can achieve
But, speedup was not shown…
GPU Technology Conference 2015
40
Reduction, etc.
Altera Stratix V GX Max bandwidth between CPUs is approx. 7GB/s with Gen3 x8,
GPU Technology Conference 2015
41
GPU Technology Conference 2015
42
Multiple orthogonal paradigms
MPI – array must be distributed and communicated (two-side or one-side) CUDA, OpenCL, OpenACC – memory allocation, data movement (to/from
controlling multiple devices if there are – CUDA 4.0 or with OpenMP
Issues
how to combine array distribution, internal-communication, external-
simple and easy-to-understand programming model is required for high
GPU Technology Conference 2015
43
PGAS language (C & Fortran) with directive base parallel
Based on our traditional PGAS language XcalableMP (XMP) OpenACC is used for control on accelerating devices Developed in AICS, RIKEN under JST-CREST joint project We implement the compiler and run-time system both for general
GPU Technology Conference 2015
44
Global View (PGAS): based on data parallel concept, directives similar to OpenMP is used for data and task distribution (easy programming)
Local View: based on local data and explicit communication (easy performance tuning)
Incremental parallelization from original sequential code
Low cost for parallelization -> high productivity
Each node processes the local data on that node
User can clearly imagine the data distribution and parallelization for easiness of tuning
Communication target of variables (arrays) and partitions can be simply specified
Communication point is specified by user, in easy manner
GPU Technology Conference 2015
45
#pragma xmp nodes p(4) declare node set #pragma xmp template t(0:99) declare template #pragma align array[i] with t(i) distribute array : owner of t(i) has a[i] #pragma xmp distribute t(BLOCK) on p distribute template
virtual array representing data(index) space
array distribution, work-sharing must be done using template
100
100
100 25 50 75
array[] 100 25 50 75 Example)
GPU Technology Conference 2015
46
GPU Technology Conference 2015
47
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
#pragma xmp gather(var=list) gather array data (collect entire elements)
GPU Technology Conference 2015
48
broadcast
barrier synchronization
reduce operation
data movement in global view
GPU Technology Conference 2015
49
GPU Technology Conference 2015
#0 #1
Distribution among nodes Distribution among ACCs.
CPU ACC
Direct Comm. between ACCs Comm. between CPUs #pragma acc device d = nvidia(0:3) #pragma xmp reflect_init (a) device #pragma xmp loop (i) on t(i) for (int i = 0; i < 100; i++){ #pragma acc kernels loop on_device(d) for (int j = 0; j < 100; j++){ a[i][j] = ... } } #pragma xmp reflect_do (a)
50
for MPI directive-base communication on distributed arrays are
using TCA for direct GPU-memory copy
GPU Technology Conference 2015
51
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE]; #pragma xmp nodes p(x, y) #pragma xmp template t(0:YSIZE−1, 0:XSIZE−1) #pragma xmp distribute t(block, block) onto p #pragma xmp align [j][i] with t(i,j) :: u, uu #pragma xmp shadow uu[1:1][1:1] … #pragma acc data copy(u) copyin(uu) { for(k=0; k<MAX_ITER; k++){ #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) uu[x][y] = u[x][y]; #pragma xmp reflect (uu) acc #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) u[x][y] = (uu[x-1][y]+uu[x+1][y]+ uu[x][y-1]+uu[x][y+1])/4.0; } // end k } // end data
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE]; #pragma xmp nodes p(x, y) #pragma xmp template t(0:YSIZE−1, 0:XSIZE−1) #pragma xmp distribute t(block, block) onto p #pragma xmp align [j][i] with t(i,j) :: u, uu #pragma xmp shadow uu[1:1][1:1] … #pragma acc data copy(u) copyin(uu) { for(k=0; k<MAX_ITER; k++){ #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) uu[x][y] = u[x][y]; #pragma xmp reflect (uu) #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) u[x][y] = (uu[x-1][y]+uu[x+1][y]+ uu[x][y-1]+uu[x][y+1])/4.0; } // end k } // end data
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE]; #pragma xmp nodes p(x, y) #pragma xmp template t(0:YSIZE−1, 0:XSIZE−1) #pragma xmp distribute t(block, block) onto p #pragma xmp align [j][i] with t(i,j) :: u, uu #pragma xmp shadow uu[1:1][1:1] … #pragma acc data copy(u) copyin(uu) { for(k=0; k<MAX_ITER; k++){ #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) uu[x][y] = u[x][y]; #pragma xmp reflect (uu) acc #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) u[x][y] = (uu[x-1][y]+uu[x+1][y]+ uu[x][y-1]+uu[x][y+1])/4.0; } // end k } // end data
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE]; #pragma xmp nodes p(x, y) #pragma xmp template t(0:YSIZE−1, 0:XSIZE−1) #pragma xmp distribute t(block, block) onto p #pragma xmp align [j][i] with t(i,j) :: u, uu #pragma xmp shadow uu[1:1][1:1] … #pragma acc data copy(u) copyin(uu) { for(k=0; k<MAX_ITER; k++){ #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) uu[x][y] = u[x][y]; #pragma xmp reflect (uu) acc #pragma xmp loop (y,x) on t(y,x) #pragma acc parallel loop collapse(2) for(x=1; x<XSIZE-1; x++) for(y=1; y<YSIZE-1; y++) u[x][y] = (uu[x-1][y]+uu[x+1][y]+ uu[x][y-1]+uu[x][y+1])/4.0; } // end k } // end data
XACC (TCA) OpenACC+MPI (GDR)
Number of Nodes Performance (GFlops)
XACC (TCA) OpenACC+MPI (GDR)
Number of Nodes
For size L, size of sleeve area is approximately 520KB, so TCA’s advantage is small compared to MVAPICH2-GDR. Additionally, TCA requires a barrier synch. after DMA transfer to cause additional overhead
better
TCA is a basic research on the possibility on direct network
Toward strong-scaling on post-peta to exascale HPC research, such
Language/Programming is also very important issue for high
XcalableACC + TCA is a solution Awarded in HPC Challenge Class2 Best Performance Award at
GPU Technology Conference 2015
57
[AsHES2015] Kazuya Matsumoto, Toshihiro Hanawa, Yuetsu Kodama, Hisafumi Fujii, Taisuke Boku, ” Implementation of CG Method on GPU Cluster with Proprietary Interconnect TCA for GPU Direct Communication,” The Internatonal Workshop on Accelerators and Hybrid Exascale Systems (AsHES2015) , May 2015 (To appear)
[CANDAR2014] Takuya Kuhara, Takahiro Kaneda, Toshihiro ihiro Hanawa wa, Yuetsu Kodama, Taisuke Boku, and Hideharu Amano, “A preliminarily evaluation of PEACH3: a switching hub for tightly coupled accelerators,” 2nd International Workshop on Computer Systems and Architectures (CSA‘14), in conjunction with the 2nd International Symposium on Computing and Networking (CANDAR 2014), pp. 377 - 381, Dec. 2014.
[WACCPD2014] Masahiro Nakao, Hitoshi Murai, Takenori Shimosaka, Akihiro Tabuchi, Toshihiro ihiro Hanawa wa, Yuetsu Kodama, Taisuke Boku, Mitsuhisa Sato, "XcalableACC: Extension of XcalableMP PGAS Language using OpenACC for Accelerator Clusters," Workshop on accelerator programming using directives (WACCPD 2014), in conjunction with SC14, pp. 27-36, Nov. 2014
[HeteroPar2014] Norihisa Fujita, Hisafumi Fujii, Toshihiro ihiro Hanawa, Yuetsu Kodama, Taisuke Boku, Yoshinobu Kuramashi, and Mike Clark, "QCD Library for GPU Cluster with Proprietary Interconnect for GPU Direct Communication," 12th International Workshop Algorithms, Models and Tools for Parallel Computing on Heterogeneous Platforms (HeteroPar2014), LNCS 8805, pp. 251- 262, Aug. 2014.
[HEART2014] Yuetsu Kodama, Toshihiro ihiro Hanawa, Taisuke Boku and Mitsuhisa Sato, "PEACH2: FPGA based PCIe network device for Tightly Coupled Accelerators," Fifth International Symposium on Highly-Efficient Accelerators and Reconfigurable Technologies (HEART2014), pp. 3-8, Jun. 2014
[HOTI2 I2013] 013] Toshihir iro
wa, Yuetsu Kodama, Taisuke Boku, and Mitsuhisa Sato, "Interconnect for Tightly Coupled Accelerators Architecture," IEEE 21st Annual Sympsium on High-Performance Interconnects (HOT Interconnects 21), short paper, pp. 79-82, Aug. 2013
[AsHE HES20 2013] 13] Toshihiro ihiro Hanawa wa, Yuetsu Kodama, Taisuke Boku, and Mitsuhisa Sato, "Tightly Coupled Accelerators Architecture for Minimizing Communication Latency among Accelerators," The Third International Workshop on Accelerators and Hybrid Exascale Systems (AsHES2013), pp. 1030-1039, May 2013.
GPU Technology Conference 2015
58
Toshihiro Hanawa
Taisuke Boku
GPU Technology Conference 2015
59