8th Symposium

学際大規模情報基盤共同利用・共同研究拠点公募型共同研究 平成２８年度採択課題

jh160014-ISH

Mohamed Wahib, RIKEN Advanced Institute for Computational Science

Data Locality Optimization Strategies for AMR Applications on GPUaccelerated Supercomputers
Problem Statement

Data-centric AMR

Motivation

GPUCPU

Adaptive Mesh Refinement (AMR), which is a model for adapting theresolution of a stencil mesh locally. AMR is one of the paths to multi-scaleexascale applications. However, producing efficient AMR code is hard,
especially for GPUs. As a result, typical AMR frameworks require the user towrite his own optimized code for the target architecture. In addition thegenerated AMR code is not optimized for reducing data transfer in GPUaccelerated supercomputers.

Copy Initial Arrays

1. Initialize

3. Compute Stencil

InvokeLoopInvoke

Octree(AMR Metadata)

Decision Kernel

Copy δInvoke<δRefine>δ

Consolidate

GPU1 Memory

Correction Kernel

4. Post-Stencil (Correction)
5. Evaluate Error

Octants(Data Arrays)

Stencil Kernel

InvokeOur approach to achieve target

Octants(Data Arrays)

Refine Kernel

InvokeGPU2 Memory

Consolidate Kernel

6. Update & Balance Octree- A compiler-based framework for producing efficient AMR code (for GPUs)- Architecture-independent interface provided to the user
- A performance model for quantifying the efficiency

7. FinalizeOctants(Data Arrays)

Copy Final Arrays

Figure 4: Conceptual overview of the proposed model. All operations touching the data arrays
are done by GPU kernels. Accordingly, CPU specializes in operations applied on the octree
while GPU specializes in operations applied on the data arrays.

Key results

Framework implementation

Our framework generates code comparable in speedup and scalability tohand-written optimized GPU AMR implementations using up to ~1000 GPUs.

The implementation is based on LLVM compiler infrastructure.
Fixed Mesh Code(Annotated)

LLVM-IRCompiler
Front EndPasses

Daino Runtime

AMR Framework

Optimized LLVM-IR

AMR library

Octree-based AMR

Call

LLVM

Comm. library

Object Files
Linker

The mesh is organized into a hierarchy of refinement levels. The mesh is
usually decomposed into relatively small fixed-sized octants of mesh cells.

Adapted Mesh Executable

K1

IRIRClangIRPassFrontEnd

C/C++

PassK1

PassK0Copy Ghost Layers

2. Exchange Ghost Layers

CPU Memory

GPU0 MemoryMachine
CodeBackEndIRLLVM proper

Figure 5: Apply translations and optimizations as compiler passes

K0Results
P2Applications

Hydrodynamics Solver: We model a hydrodynamics application using Euler
equations extending the GAMER implementation [2].
Shallow-water Solver: We model shallow water simulations by depthaveraging the Navier–Stokes equations.
Phase-field Simulation: We evaluate an AMR version of a phase-field
simulation for modeling 3D dendritic growth [3].
2.5E+03
2.0E+03

PHASE-FIELD

Uniform MeshAuto AMR

2.5E+02

Runtime (Seconds)

Framework

5.0E+02

5.0E+02
0.0E+00

48163264128256512

996

0.0E+00

4

Number GPUs (Mesh size per GPU: 4,096x512x512)

8163264128

2565123996

2.0E+02
1.5E+02
1.0E+02
5.1E+01
1.0E+00

48163264128256512

Number GPUs (Mesh size per GPU: 8,1923)

Number GPUs (Mesh size per GPU: 4,096 )996

Figure 6: Weak scaling of uniform mesh, hand-written and automated AMR (GAMER-generated AMR included in hydrodynamic)

CPU AMR Executable
GPU AMR ExecutablePHASE-FIELD

CPU AMR Executable

Uniform Mesh
Auto AMR
Hand-written AMR

3.5E+04
3.0E+04

Architecture-independent interface
The framework is composed ofa compiler and runtime. The input to the framework is serial
code applying stencils operations on a uniform grid. The useradds directives to identify thestencil functions and relevant
data arrays.

9.6 x

2.0E+02

2.0E+04

1.5E+02

1.5E+04
1.0E+04

1.5 x

1.0E+02

8.9 x

5.1E+01

5.0E+03
1.0E+00

3.0E+02

SHALLOW-WATER S

Uniform Mesh
Auto AMR (Daino)
Hand-written AMR
Auto AMR (GAMER)

3.5E+02

2.5E+02

2.5E+04

Runtime (Seconds)

Figure 2: An illustration of the architecture-neutral interface used in the framework

1.7 x

HYDRODYNAMICS

2.1E+03

Runtime (Seconds)

GPU AMR Executable

3.0E+02

1.0E+03

1.0E+03

Framework

Uniform Mesh

Auto AMR (GAMER)

2.0E+03

1.5E+03

1.5E+03

SHALLOW-WATERS
Auto AMR
Hand-written AMR

HYDRODYNAMICS
Auto AMR (Daino)

Uniform Mesh
Hand-written AMR

Hand-written AMR

3.1 x#pragma daino kernel
void 3D_alloy(..){#pragma daino data (Nx,Ny,Nz){p, u, dpt, no, o;}… kernel code ...
}

4.1 x

void 3D_alloy(..){#pragma omp for… kernel code ...
}Runtime (Seconds)__global__ 3D_alloy(..){… CUDA kernel code ...
}Uniform Mesh Serial C Code48163264

128256512

Number GPUs (Mesh size 4,0963)

996

1.0E+00

48163264128

256512996

Uniform Mesh
Auto AMR
Hand-written AMR

3.0E+02
2.5E+02

Runtime (Seconds)

OpenMP Code

Runtime (Seconds)

CUDA code

8.7 x

Our framework

Other AMR frameworks

9.6 x

Figure 1: An example of a quadtree for a 2D AMR mesh (work divided on three processors)

1.7 x

P1

1.9 x

P0

2.0E+02

4.1 x

1.5E+02
1.0E+02

3.2 x
5.1E+01
1.0E+00

Number GPUs (Mesh size per GPU: 4,0963)

48

1632

64

128

256512

Number GPUs (Mesh size per GPU: 8,1923)

996

Figure 7: Strong scaling of uniform mesh, hand-written and automated AMR (GAMER-generated AMR included in hydrodynamic)

Table 1: Phase-field runtime breakdown (%) for Daino AMR

Table 2: Hydrodyn. runtime breakdown (%) for Daino AMR

Table 3: Shallow. runtime breakdown (%) for Daino AMR

Figure 3: Minimal example directives

学際大規模情報基盤共同利用・共同研究拠点 第８回シンポジウム
2016年 7月 14日，15日

THE GRAND HALL （品川）