XcalableMP and XcalableACC
for Productivity and Performancein HPC Challenge Award Competition†‡††

Masahiro Nakao, Hitoshi Murai, Hidetoshi Iwashita† Akihiro Tabuchi,* Taisuke Boku‡*
Takenori Shimosaka,†‡*Mitsuhisa Sato

† RIKEN Advanced Institute for Computational Science, Japan‡ Center for Computational Sciences, University of Tsukuba* Graduate School of Systems and Information Engineering, University of Tsukuba
HPC Challenge Class II BoF@SC14, Nov. 18th

Outline
1. XcalableMP (XMP) for cluster systems(14min.)2. XcalableACC (XACC) for accelerator cluster systems(6min.)

Extension of XMP using OpenACCSorry !!, work-in-progress
The submission report is available at http://xcalablemp.org

2

What is XcalableMP (XMP) ?
Directive-based language extensions of Fortran and C
By XMP speciﬁcation working group of PC cluster consortium(SC Booth#2924)
Version 1.2.1 speciﬁcation available (http://xcalablemp.org)

Support two memory models
Global-view (HPF-like data/work mapping directives)
Local-view (coarray)

Implementation of Compiler
Omni XMP Compiler version 0.9 (http://omni-compiler.org)
Platforms: Fujitsu the K computer and FX10, Cray XT/XE,
IBM BlueGene, NEC SX, Hitachi SR, Linux clusters, etc.
3

Code example (Global-view)
int a[MAX];#pragma xmp nodes p(4)#pragma xmp template t(0:MAX-1)#pragma xmp distribute t(block) on p
#pragma xmp align a[i] with t(i)main(){
int i, j, res = 0;

Data distribution

add to the serial code : incremental parallelization#pragma xmp loop on t(i) reduction(+:res)for(i = 0; i <MAX; i++){Work mapping and data synchronizationa[i] = func(i);res += array[i];}4Code example (Local-view)
double a[100]:[*], b[100]:[*];
int me = xmp̲node̲num();

Deﬁne Coarrays

if(me == 2)
a[:]:[1] = b[:];Put Operationif(me == 1)a[0:50] = b[0:50]:[2];Get Operation

Coarray synax in XMP/C

array_name[start:length]:[node_number];
XMP/Fortran is upward compatible with Fortran 2008
5Results and Machine
Summary

Four HPCC Benchmarks

Benchmark# Nodes

Performance (/peak)

SLOC

Ver. 1

16,384

971 TFlops (46.3%)

313

Ver. 2

4,096

423 TFlops (80.7%)

426

FFT

82,944212 TFlops (2.0%)

205

STREAM

82,994

3,583 TB/s (67.5%)

69

RandomAccess

16,384254 GUPs

253HPLThe K computer: 82,944 nodes
SPARC64 VIIIfx Chip, 128 GFlops
DDR3 SDRAM 16GB, 64GB/s
Tofu Interconnect
6D mesh/torus network
http://www.aics.riken.jp/jp/outreach/photogallery.html

5GB/s x 4links x 26

HPL version 1
Source lines of Code (SLOC) is 313, written in XMP/CBlock-Cyclic Distribution

A[N][N]

double	  A[N][N];	  #pragma	  xmp	  nodes	  p(P,Q)	  #pragma	  xmp	  template	  t(0:N-­‐1,	  0:N-­‐1)	  #pragma	  xmp	  distribute	  t(cyclic(NB),	  \	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  cyclic(NB))	  onto	  p	  #pragma	  xmp	  align	  A[i][j]	  with	  t(j,i)

1

2

3

4

1

2

Programmer can use BLAS for distributed array.

Panel Broadcast by using gmove directive

A[N][N]

A_L[N][NB]double	  A_L[N][NB];	  #pragma	  xmp	  align	  A_L[i][*]	  with	  t(*,i)	  k	  	  	  	  :	  #pragma	  xmp	  gmove	  A_L[k:len][0:NB]	  =	  A[k:len][j:NB];
lenjNB7

HPL version 2
SLOC is 426, written in XMP/C”Lookahead algorithm” by using gmove directive with async clause
Overlap communication and calculation
double	  A_L[N][NB];	  #pragma	  xmp	  align	  A_L[i][*]	  with	  t(*,i)	  	  	  	  	  :	  k#pragma	  xmp	  gmove	  async(1)	  A_L[k:len][0:NB]	  =	  A[k:len][j:NB];	  	  	  	  	  :	  lenfor(m=j+NB;m<N;m+=NB){	  	  	  for(n=j+NB;n<N;n+=NB){	  	  	  	  	  cblas_dgemm(&A[m][n],	  ..);	  	  	  	  	  if(xmp_test_async(1)){	  	  	  	  	  	  	  	  //	  receive	  A[k:len][j:NB];	  	  	  	  	  	  	  	  	  	  	  	  :A[N][N]

jA_L[N][NB]

NB

asynchronous broadcast
communication

Confirm whether data with async clause comes or not.
8

Performance of HPL
1000#TFlops

109 TFlops (83.5%)
1,024 nodes
100#

10#
256#

1niosrVe

423 TFlops (80.7%)
4,096 nodes

2noisreV

971 TFlops (46.3%)
16,384 nodes

310 TFlops (59.1%)
4,096 nodes

88 TFlops (67.2%)
1,024 nodes
1024#

4096#

16384#

Number of nodes

XMP-HPL Version 2 has a good scalability.
Sorry, the measurement in 16,384 nodes is late for this BoF.
9RandomAccess
SLOC is 253, written in XMP/C
Local-view programming with XMP/C coarray syntax
The XMP RandomAccess is iterated over sets of CHUNK updates oneach nodeu64Int	  recv[LOGPROCS][RCHUNK+1]:[∗];	  ...	  
for	  (j	  =	  0;	  j	  <	  logNumProcs;	  j++)	  {	  	  	  recv[j][0:num]:[i_partner]	  =	  send[i][0:num];	  	  Deﬁne coarrayPut operation#pragma	  xmp	  sync_memory	  #pragma	  xmp	  post(p(i_partner),	  0)	  	  	  	  :	  #pragma	  xmp	  wait(p(j_partner))	  }A point-to-point synchronization is speciﬁed with the XMPʼs post and wait
directives to realize asynchronous behavior of this algorithm
10Performance of RandomAccess
Last year, to implement the post/wait directives, XMP uses MPI̲Send/Recv.
This year, to implement them, XMP uses RDMA of the K computer.

1000"

254 GUPs
16,384 nodes

raeysihT

GUPs

100"

Las

10"

raety

162 GUPs
16,384 nodes

1"
64"

256"

1024"

4096"

16384"

Number of nodes
11FFT and STREAM
Code cleanup and performance improvement.
Please refer to the submission report at http://xcalablemp.org
FFT (SLOC 205, XMP/F)

Th
isTFlops

100"

Las

10"

50 TFlops
38,864 nodesraety

3,583 TB/s
82,944 nodes

10000#

1000#si
ThTB/syear

212 TFlops
82,944 nodes

STREAM (SLOC 69, XMP/C)

La

100#

1"
512"

4096"

32768"

Number of nodes

10#
1024#

rayeyts

8192#

rae

2,439 TB/s
82,944 nodes

65536#

Number of nodes
12Compare to two versions
Improvement rate (on the same nodes) 37 - 94% improvement !!
1.94

Ratio

2.0##
1.5##

1.37

1.56

1.47

1.0##

Good

0.5##
0.0##
HPL#

(4,096 nodes)

RandomAccess#

STREAM#(16,384 nodes) (16,384 nodes)

FFT#(36,864 nodes)

SLOC
2500"

Last year, work-in-progress to clean up code

2416

2000"
1500"
1000"
500"

313 426

250 253

HPL"

RandomAccess"

0"

66

69

STREAM"

205

GoodFFT"

13Outline
1. XcalableMP (XMP) for cluster systems(14min.)2. XcalableACC (XACC) for accelerator cluster systems(6min.)

Extension of XMP using OpenACC
Sorry !!, work-in-progress
The submission report is available at http://xcalablemp.org

14

What is XcalableACC?
Extension of XMP using OpenACC for accelerator clusters
Feature:Mix XMP and OpenACC directives seamlessly
Support transferring data among accelerators directly

15

Diﬀerence XMP and XACC memory models
XMP memory model
Global Indexing

・・
node #1

Transfer data amongnode #2

Host memories (XMP)・・
HostHostXACC memory model

Map “global Indexing” to accelerators

Global Indexing・・
node #1

Transfer data among
Host memories (XMP)node #2

HostACC

HostACC・・

Transfer data among
Host - ACC (OpenACC)
Transfer data among
ACCs (XACC)

16

XACC code example
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];	  Laplaceʼs equation#pragma	  xmp	  reflect	  (uu)	  acc	  Exchange halo region of 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

Data Distribution
Transfer XMP distributed arraysto accelerator
OpenACC directive parallelizes
the loop statement parallelized
by XMP directive

When “acc” clause is speciﬁed in
XMP communication directive,
data on accelerator is transferred.

17Results and Machine
Summary
BenchmarkThree HPCC Benchmarks and HIMENO Benchmark#Nodes #CPUs #GPUs Performance (/peak)

SLOC

HPL

32

64

128

7 TFlops (4.2%)

343

FFT

3264

-257 GFlops (0.1%)

205

STREAM

64

128

256

15 TB/s (20.4%)

84

HIMENO

64

128

256

14 TFlops (1.4%)

253

HA-PACS/TCA: 64 nodes
Ivy Bridge E5-2680v2, 224GFlops x 2 Sockets
DDR3 SDRAM 128GB, 59.7GB/s x 2
Inﬁniband 4xQDR x 2 rails : 8GB/s
NVIDIA K20X (4GPUs / Node)1.31 TFlops/GPU(SP), 3.95 TFlops/GPU(DP)250GB/s/GPU
http://www.ccs.tsukuba.ac.jp/CCS/eng/research-activities/projects/ha-pacs

18

STREAMThe XACC STREAM uses both CPUs and GPUs together,XMP, OpenACC, and OpenMP directives are used.
#pragma xmp nodes p(*)#pragma acc data copy(a[0:GPU_SIZE], b[0:GPU_SIZE], c[0:GPU_SIZE]){
for(k=0; k<NTIMES; k++) {#pragma xmp barrier
times[k] = -xmp_wtime();#pragma acc parallel loop async
for (j=0; j<GPU_SIZE; j++)a[j] = b[j] + scalar*c[j];on GPU#pragma omp parallel forfor (j=GPU_SIZE; j<MAX_SIZE; j++)a[j] = b[j] + scalar*c[j];on CPU#pragma acc wait

Wait until GPU task completes#pragma xmp barrier
times[k] += xmp_wtime();}} // acc data

19

Performance of STREAM
reasonable performance
100000PUC(CCAX
84:CSLO

GB/s10000)
UP+G

5 TB/s64 nodes1000UPCnlyO(XMP : 69COLS10015 TB/s
64 nodes (256GPUs))

10
12

4

8

1632

64

Number of nodes
20

HIMENO Benchmark
Stencil application of incompressible ﬂuid analysis code
Solving the Poissonʼs equation
Sequential and MPI Version HIMENO Benchmark is available at
http://accc.riken.jp/2444.htm
float	  p[MIMAX][MJMAX][MKMAX];	  //	  Define	  distributed	  array	  and	  halo	  #pragma	  acc	  data	  copy(p)	  ..	  
{	  ..	  
#pragma	  xmp	  reflect	  (p)	  acc	  ..	  	  #pragma	  xmp	  loop	  (k,j,i)	  on	  t(k,j,i)	  #pragma	  acc	  parallel	  loop	  ..	  
for(i=1;	  i<MIMAX;	  ++i)	  	  	  for(j=1;	  j<MJMAX;	  ++j){	  #pragma	  acc	  loop	  vector	  ..	  	  	  	  	  for(k=1;	  k<MKMAX;	  ++k){	  	  	  	  	  	  	  S0	  =	  p[i+1][j][k]	  *	  ..;

Transfer distributed array toaccelerator
Exchange halo region
Parallelize loop statement
Only add XMP and OpenACC
directives into the sequential
Himeno benchmark.
21

Performance of HIMENO
reasonable performance100000)
UPGyln(oCC
3XA
12:CSLO

GFlops10000

1000NEM
IHnoisrev
MPI : 325COL
S

100

10124816

14 TFlops
64 nodes (256GPUs))
UPCylnO (O32

1.6 TFlops
64 nodes64

Number of nodes
22

HPL and FFT
Sorry !! work-in-progress for implementing and tuning.
FFTHPLU)

PG

ylno( 34C
3C C:AX LO
S
7 TFlops (4.2%)
32 nodes(128GPUs)

1000

using cuBLAS
100
12481632Cyln(o 05P :2M
X OC
SL100GFlops

GFlops

10000)
UP

257 GFlops (0.1%)
32 nodes10will use FFTE-CUDA
1
1

Number of nodes

24816

32

Number of nodes

Time of transfer data between CPU andhost memory dominates the total computation time23

Conclusion
XMP on the K computer
Benchmark# Nodes

Good productivity andperformance !!
Performance (/peak)

SLOC

Ver. 1

16,384

971 TFlops (46.3%)

313

Ver. 2

4,096

423 TFlops (80.7%)

426

FFT

82,944212 TFlops (2.0%)

205

STREAM

82,994

3,583 TB/s (67.5%)

69

RandomAccess

16,384254 GUPs253

HPL

XACC on HA-PACS/TCA
BenchmarkWe will improve HPL and FFTnext year.

#Nodes #CPUs #GPUs Performance (/peak)

SLOCHPL

32

64

128

7 TFlops (4.2%)

343

FFT3264-257 GFlops (0.1%)

205

STREAM

64

128

256

15 TB/s (20.4%)

84

HIMENO64

128

256

14 TFlops (1.4%)

253
24

For more information
Please visit our booth !!
RIKEN AICS (Advanced Institute for Computational Science) #2413
Center for Computational Sciences, University of Tsukuba #3215

25