XcalableACC: Extension ofXcalableMP PGAS Language using
OpenACC for Accelerator Clusters
Masahiro Nakao,† Hitoshi Murai,† Takenori Shimosaka†Akihiro Tabuchi,‡ Toshihiro Hanawa, Yuetsu Kodama‡ *†‡ *Taisuke Boku,‡* Mitsuhisa Sato
† RIKEN Advanced Institute for Computational Science, Japan‡ Graduate School of Systems and Information Engineering, University of Tsukuba
Information Technology Center, The University of Tokyo* Center for Computational Sciences, University of Tsukuba

1st WACCPD, November 17th, 2014

BackgroundAccelerator cluster
High performanceHigh eﬃciency power consumption(June, 2014)
Tsubame KFC, Tokyo Techhttp://www.titech.ac.jp

HA-PACS, Univ. of Tsukuba(Our Cluster)
http://www.ccs.tsukuba.ac.jp

2

Problems (1/2)
1. Complex programming on accelerator clusters
OpenACC and Message Passing Interface (MPI)MPI_Init(&argc,	  &argv);	  MPI_Comm_rank(MPI_COMM_WORLD,	  &rank);	  double	  a[100];	  #pragma	  acc	  data	  copyin(a)	  {	  #pragma	  acc	  parallel	  loop	  	  	  for(i=0;i<100;i++)	  	  	  	  	  a[i]	  =	  i;	  }	  if(rank	  ==	  0)	  	  	  MPI_Send(a,	  100,	  MPI_DOUBLE,	  1,	  tag,	  	  	  	  	  	  	  	  	  	  	  MPI_COMM_WORLD);	  else	  	  	  MPI_Recv(a,	  100,	  MPI_DOUBLE,	  0,	  tag,	  	  	  	  	  	  	  	  	  	  	  MPI_COMM_WORLD,	  &status);

OpenACC has a good productivity.
But MPI is often diﬃcult.
Must control local data byusing primitive MPI functions(While search arguments ...)

3

Solutions (1/2)
1. Usage of XcalableMP (XMP)
Directive-based language extensions of Fortran and C language
for clusters
int	  a[MAX];	  MPIMPI_Init(&argc,	  &argv);	  MPI_Comm_rank(MPI_COMM_WORLD,	  &rank);	  MPI_Comm_size(MPI_COMM_WORLD,	  &size);	  dx=MAX/size;	  llimit=rank*dx;	  if(rank	  !=(size-­‐1))	  ulimit=llimit*dx;	  else	  ulimit	  =	  MAX;	  temp_res=0;	  for(i=llimit;i<ulimit;i++){	  	  	  a[i]=func(i);	  	  	  tmp_res	  +=	  a[i];	  }	  MPI_Allreduce(&temp_res,	  &res,	  1,	  ..);

int	  a[MAX];	  XMP#pragma	  xmp	  nodes	  p(*)	  #pragma	  xmp	  template	  t(1:MAX)	  #pragma	  xmp	  distribute	  t(block)	  onto	  p	  #pragma	  xmp	  align	  a[i]	  with	  t(i)	  res=0;	  #pragma	  xmp	  loop	  t(i)	  reduction(+:res)	  for(i=0;i<MAX;i++){	  	  	  a[i]=func(i);	  	  	  res	  +=	  a[i];	  }XcalableACC (XACC): Extension of XMP using OpenACC
for accelerator clusters
4

Problems (2/2)
2. Large communication latency among accelerators on diﬀerent nodes(1)

ACC. = Accelerator Mem. = Host memory

(3)
ACC

ACC

CPUCPU(1) Copy data from ACC. to Mem. by OpenACC(2) Send data from Mem. to Mem. by MPI(3) Copy data from Mem. to ACC. by OpenACC

Twice Mem. copies occur !!

MemMem

NIC

NIC

(2)
SW

To overcome

(1)
ACC

ACC

CPUCPU

GPUDirect does NOT require Mem. copy.
Faster than the above method
But accelerator applications want smaller latency !!

Mem

GPUDirect still requires time for “MPI software
stack”, and “Protocol conversion (e.g. PCIe to InﬁniBand).”

Mem

NIC

NIC

SW

5

Solutions (2/2)
2. Usage of Tightly Coupled Accelerators (TCA) [1]
Communication architecture based on PCIe technology
Developed by HA-PACS Project in Univ, of Tsukuba, JapanNodes are connected using PCIe external cablethrough PEACH2, which is a TCA interface Board
Direct, low latency data transfers amongaccelerator memories

(1)
ACC

PE2

ACC

PE2

No host memory copies
No MPI software stack

CPUCPUNo protocol conversions
MemNIC

SWMem

NIC

PE2 = PEACH2[1] Toshihiro Hanawa et al. "Tightly Coupled Accelerators Architecture for Minimizing
Communication Latency among Accelerators," in IPDPSW ʼ13 Proceedings of the 2013

6

Solutions (2/2)
TCA provides Hardware-level APIs to deal with itself.
It is diﬃcult !!
TCA requires a programming interface.

We also develop XACC as a TCA programming interface.

7

Objectives(1) Provide a design for a directive-based language foraccelerator clusters and evaluates its eﬀectiveness(2) Discuss our experiences implementing directives fordata transfer among accelerators

8

Outline
Overview of XMP and XACC programming models
Implementation of XACC compiler
Evaluation of productivity and performance of XACC
programming model

9

Outline
Overview of XMP and XACC programming models
Implementation of XACC compiler
Evaluation of productivity and performance of XACC
programming model

10

XMP Global-view memory model
PGAS (Partitioned Global Address Space) is a memory region can beaccessed from other processes
In XMP, to improve productivity, programmers can use “Global indexing.”

01 2 3 45 6 7 8 9 10 11 12 13 14 15

Global indexing

Node 1

Distributed array

Node 2
Node 3
Node 4

two elements copy
#pragma	  xmp	  gmove	  a[0:2]	  =	  a[4:2];
Maintain a sequential code image, as OpenACC and OpenMP
11

XMP Global-view memory model
PGAS (Partitioned Global Address Space) is a memory region can beaccessed from other processes
In XMP, to improve productivity, programmers can use “Global indexing.”

01 2 3 4

5 6 7 8 9 10 11 12 13 14 15

Global indexing

Node 1
Node 2

halo copy amongneighbor nodes

Node 3#pragma	  xmp	  shadow	  a[1:1]	  ..	  	  Node 4#pragma	  xmp	  reflect	  (a)	  #pragma	  xmp	  loop	  on	  t(i)	  for(..){	  	  	  ..	  =	  array[i-­‐1]	  +	  array[i+1]; For Stencil Application
Maintain a sequential code image, as OpenACC and OpenMP
12

Reason for extending XMP to use OpenACC
Basically, programmers can use OpenACC directives in XMP source code.

Problems:
1. XMP does NOT support XMP distributed array in OpenACC directives
Local OpenACC directive, data movement between CPU and accelerator,
cannot be applied to XMP distributed array
2. XMP does NOT express transferring data among accelerators directly
Need to express data movement ACC. -> MEM. -> MEM. -> ACC.

Solutions: Development of XACC for accelerator clusters
1. XACC supports XMP distributed array in OpenACC directives
Programmer can use both XMP and OpenACC directives seamlessly
2. XACC can express transferring data among accelerators directly

To implement it , TCA is used as a data transfer architecture
13

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

・・
node #1

Transfer data amongnode #2

Host memories (XMP)

・・
Host

HostXACC 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)

14

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 and Halo
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.

15

Outline
Overview of XMP and XACC programming models
Implementation of XACC compiler
Evaluation of productivity and performance of XACC
programming model

16

Implementation of XACC compiler
Extension of the Omni XMP Compiler
Free software http://omni-compiler.org
source-to-source compiler
PGI, Cray, Omni OpenACC compilers
are available as a backend compiler

Development points:
1. Specify an XMP distributed array in OpenACC directives
2. Mix XMP loop directive and OpenACC loop directive
3. Transfer data among accelerators17

Implementation of XACC compiler
Extension of the Omni XMP Compiler
Free software http://omni-compiler.org
source-to-source compiler
PGI, Cray, Omni OpenACC compilers
are available as a backend compiler

Development points:
1. Specify an XMP distributed array in OpenACC directives Please refer toour paper
2. Mix XMP loop directive and OpenACC loop directive
3. Transfer data among accelerators

18

Transfer data among accelerators
How to transfer data among accelerators in the XACC compiler
1. DMA of PEACH2 : sophisticated DMA functions
Internal memory mode
Host memory mode

The maximum number ofregistrable regions is 1,024

2. GPUDirect
3. CUDA + MPI : ACC. -> MEM. -> MEM. -> ACC.

Problem:
In preliminary evaluation, depending on data size and hardware,“GPUDirect” may be faster than “DMA of PEACH2.”

Solution:
Implement of ﬂow of switching communication method19

Preliminary Evaluation
GPUDirect uses InﬁniBand 4xQDR x 2rails, the bandwidth is 8GB/s
DMA of PEACH2 uses PCIe Gen.2 x 8links, the bandwidth is 4GB/s
Transferring more than 1M byte, the performance ofGPUDirect is better than that of PEACH2 in PCIe Gen2

Latency (u second)

1000"
1000

GPUDirect RDMA(MVAPICH2-GDR 2.0b)

Protocol switch ofMVAPICH2-GDR

100"
100

Host memory mode(PEACH2)

1010"

Internal memory mode
(PEACH2)
1"
18"
864"
64512"
5124096"4K

32768"

32K

262144"
256K

better2097152"

2M

Transfer data (Byte)

20

Flow of communication method
START
Support
PEACH2 ?

No

Yes
Data < 1MB ?

No

SupportGPUDirect ?

Yes
Num<= 1024 ?

No

Yes
No

Yes
Internal Memory
Mode (PEACH2)

Host Memory
Mode (PEACH2)

MVAPICH2-GDR

MPI + CUDA

21

Outline
Overview of XMP and XACC programming models
Implementation of XACC compiler
Evaluation of productivity and performance of XACC
programming model

22

Target Benchmark
Himeno Benchmark
Stencil application of Incompressible ﬂuid analysis code
Solving the Poissonʼs equation
http://accc.riken.jp/2444.htm : Sequential/MPI versions
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 XMP distributed arrayto accelerator
Exchange halo region
Parallelize loop statement

23

Productivity
For comparison purposes, we also implemented HIMENO Benchmark using
OpenACC based on MPI HIMENO Benchmark (OpenACC+MPI HIMENO)
To parallelize loop statements, we use OpenACC data and loop directives
To use GPUDirect, we insert OpenACC host̲data directive before MPI functions

Source Lines of CodeBreakdown
XMP
OpenACC + MPI
XACC

OpenACC

TotalOthers-15473

488

289176213

50%
reduction

OpenACC+MPI HIMENO Benchmark requires numerous lines tocalculate indexes for a loop statement, and to transfer the halo region
In XACC, we only add directives to the sequential HIMENO Benchmark
XACC HIMENO Benchmark maintains a sequential source code image
24

Performance

HA-PACS, Univ. of Tsukuba

NVIDIA K20X
InfiniBand 4xQDR x 2rails
PCIe Gen2 x8 for PEACH2
MVAPICH-GDR2.0b
gcc-4.7, CUDA6.0, OmniOpenACC Compiler 0.9b

Performance (GFlops)

Array size 128x128x256 : Strong scaling, 1GPU/procs
320

XACC (PEACH2)
OpenACC+MPI (GDR)

x2.7 up

240
160
800better
124816

Number of Processes

In XACC, all communication uses the Internal memory mode of PEACH2Internal memory mode of PEACH2 vs. GPUDirect RDMA over InﬁniBand
Note that performance of OpenACC using PEACH2is the almost the same that of XACC.
25

Conclusion
New programming model for accelerator clusters, called XACC
Programmer can use both XMP and OpenACC directives seamlessly
Programmer can express transferring data among accelerators directly
Two Objectives(1) Provide a design for a directive-based language for accelerator clusters
and evaluates its eﬀectivenessThrough implementing the HIMENO Benchmark,
XACC has a good productivity and good performance using TCA(2) Discuss our experiences implementing directives for data transfer
among accelerators via TCA

Automatically select appropriate communication method

26

Future works
Development of mini-applications and real world applications
Usage of both TCA and normal interconnect (e.g. InﬁniBand)
Extension of OpenACC to use multiple accelerators
Need to specify a device number by acc̲set̲device̲num()
before each statement to use multiple accelerators
To use multiple accelerators seamlessly similar to XMP directives
float	  a[N],	  b[N];	  #pragma	  acc	  device	  d(*)	  #pragma	  acc	  declare	  device_resident(a)	  layout([block])	  ¥	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  shadow([1:1])	  on_device(d)	  	  ..	  
#pragma	  acc	  reflect	  (a)	  #pragma	  acc	  parallel	  loop	  layout(a[i])	  on_device(d)	  for	  (int	  j	  =	  1;	  j	  <	  99;	  j++)	  	  	  b[i]	  =	  a[i-­‐1]	  +	  a[i+1];

Also show
SC poster27