XcalableACC: Extension of XcalableMP using OpenACC for
Accelerator Clusters
Motivation

Example 1: Loop parallelization

In order to develop an application on accelerator clusters, adouble a[N];#pragma xmp nodes p(*)
Define XMP distributed array#pragma xmp template(0:N-1)#pragma xmp distribute t(block) onto p#pragma xmp align a[i] with t(i)Transfer XMP distributed array#pragma acc data copy(a)to accelerator{#pragma xmp loop on t(i)OpenACC parallel loop directive#pragma acc parallel loop
parallelizes a loop statementfor(int i=0;i<N;i++){a[i] = .. ;parallelized by XMP loop directive

programmer selects a combination of MPI and one of the available
accelerator programming models, such as OpenACC, OpenCL orCUDA. In particular, OpenACC provides good productivity and highperformance. However, MPI programming is often difficult because
programmers need to distribute, transfer, and reduce data by using
primitive MPI functions.

XcalableACC
Based on studies of XMP and OpenACC, we believe that aExample 2: Usage of multiple accelerators
Define multiple accelerators#pragma acc device d(*)
float a[N], b[N];#pragma acc declare device_resident(a,b) ¥Declare distributedlayout([block]) shadow([1:1]) ¥arrays with a halo
on_device(d)..
#pragma acc reflect (b)Synchronize a halo region#pragma acc parallel loop layout(a[i]) ¥on_device(d)Parallelize a loopfor (int i=1; i<N-1; i++){statementa[i] = b[i-1] + b[i+1];..

directive-based language is effective for also accelerator clusters.
Thus, we propose a new programming model called XcalableACC(XACC), which is an XMP extension that uses OpenACC. In XACC,programmers can use both XMP and OpenACC directives
seamlessly. Therefore, XACC allows programmers to easily develop
applications on accelerator clusters. Additionally, XACC also
provides directives to transfer data among accelerators efficiently.
“Example 1” on the right side shows an example of XACC.

Evaluation of the HIMENO benchmark

Extension of OpenACC for multiple accelerators

The HIMENO Benchmark is a stencil application benchmark thatIn the OpenACC specification 2.0a, it is difficult for a single processevaluates the performance of incompressible fluid analysis code.

to use multiple accelerators. For example, programmers must specify

float p[MIMAX][MJMAX][MKMAX];#pragma xmp shadow p[1:1][1:1][0]..
#pragma xmp reflect (p) acc..
#pragma xmp loop (k,j,i) on t(k,j,i)#pragma acc parallel loop collapse(3) ..
for(i=1 ; i<MIMAX ; ++i)
for(j=1 ; j<MJMAX ; ++j){
for(k=1 ; k<MKMAX ; ++k){S0 = p[i+1][j][k] * ..;an accelerator number using the acc_set_device_num() before
each statement, calculate indexes for loop statement on each GPU,and move data among accelerators manually. To reduce theprogramming cost, we consider an extension of OpenACC to use
multiple accelerators. We believe that this extension is also useful for
XACC. “Example 2” on the right side shows an example of theextension of OpenACC.

Define a halo region
Synchronize a halo region
Parallelize a loop statement

Productivity

Definition and synchronization of a halo region

Total

To develop stencil applications easily, XACC provides shadow andDirective
XcalableMP

OpenACC

Except fordirectives

reflect directives. These directives can be used to define a halo

XcalableACC

204229173

region for a distributed array and synchronize it among nodes. InMPI + OpenACC

325-15310Omni compiler supports Tightly Coupled Accelerators(TCA) inaddition to the combination of CUDA and MPI, and GPUDirect
RDMA over InfiniBand. TCA is a communication architecture basedon PCI Express (PCIe) technology. Computation hosts with TCA
network are connected with each other using PCIe external cables,
which permits direct, low latency data transfers among accelerator
memories because they do not need host memory copies, a protocol
conversion, and MPI software stack. For more details of TCA,
please go to “Center for
Computational Sciences,
University of Tsukuba”(#3215 booth).

The source lines of the XACC HIMENO Benchmark is less than that of
the MPI+OpenACC HIMENO Benchmark. Especially, to develop theXACC HIMENO benchmark, a programmer only adds XMP andOpenACC directives into the sequential Himeno Benchmark.

Performance
Performance (GFlops)

order to transfer data among accelerator memories efficiently, the400300

The performance of XACC

XcalableACC

using TCA is up to 2.7 times

MPI + OpenACC

better than that of MPI + Open200

ACC using MVAPICH2-GDR

10002.0b on HA-PACS/TCA
1248
Number of Nodes

16system (NVIDIA K20X,
InfiniBand QDR x2 rails).