Lift: a Functional Approach to Generating High
Performance GPU Code using Rewrite Rules

ToomasRemmelg

MichelSteuwer

ChristopheDubach

The 4th South of England Regional Programming Language Seminar
27th September 2016

Heterogeneity Everywhere

GPUMobile
devices

Gameconsoles

FPGAPCs

CPU/GPU

Data centres
Supercomputers

Top 500 with parallel accelerators

Top 500 with parallel accelerators

increasing

Top 500 with parallel accelerators
new ones appearing
regularly

Top 500 with parallel accelerators

Difficult to program
Difficult to achieve high performance
Moving target

But also domain-specificaccelerators►

E.g. Google neural network accelerators•►

TPU (Tensor Processor Units)

E.g. Movidius vision accelerators•VPU (Vision Processing Units)Current situation
Applications
Device-specificinterface

OpenCL

ParallelHardwareGPUs

OpenMPVHDL

C library...

Multicore CPUs

FPGAs

NNAccelerator...

Programmers have to beexpert in high performance computing!

+ Performance is not portable!
►hand-written implementations for each device•

parametric auto-tuner

+ Performance is not portableeven on same device class
Matrix-matrix multiplication auto-tuner

+ Performance is not portableeven on same device class
Matrix-matrix multiplication auto-tuner
best
implementation

+ Performance is not portableeven on same device class
Matrix-matrix multiplication auto-tuner
best
implementation

+ Performance is not portableeven on same device class
Matrix-matrix multiplication auto-tuner
best
implementation

Auto-tuning alone fails to achieve portable performance

Need for high-level programming+Code generation

High-Level programming approaches
for heterogeneous devices
►

►

►

►

►

StreamIt (MIT)•Dataflow programming, architecture Independent•

Multiple backend (tile architecture, C, FPGA, GPU)

LiquidMetal (IBM)•Java + dataflow programming + (map / reduce)•

C, OpenCL and FPGA backend

Green-Marl (Stanford / Oracle)•

DSL for graph analysis•

operations described on graph's node or edgesl

Halide (MIT)•DSL for imagine processing, functional in nature•

GPU code generator

TensorFlow (Google)•

DSL library for AI applications•

data flow modelData flow & functional approaches:►

Hardware complexity is hidden away•►

→ code portability

Empowers the compiler•

no need for complicated analysis•

explicit data movement, no global state•

can be mapped on various hardware•scheduling is easier

General purpose framework fordata parallelism►►Delite (Stanford + EPFL)•framework for developing high-performance DSLs•

provide built-in parallel primitives•

Multicore + GPU code generation

Lift (Edinburgh)•intermediate data-parallel language•

compiler optimisations expressed as rewrite rules•

Multicore + GPU code generation•+ more to come (MPI, OpenMP, VHDL)What we have
Applications
Device-specificinterface

OpenCL

ParallelHardwareGPUs

OpenMPVHDL

C library...

Multicore CPUs

FPGAs

NNAccelerator...

What we need
ApplicationsOpenMP

VHDL

C library...

Multicore CPUs

FPGAs

NN
Accelerator...

OpenCLGPUs

What we need
Applications
Domain Specific
Languages (DSLs)

Image
Processing

NeuralNetworks

GraphAnalytics

Language forParallelism

Data-Parallel
Intermediate Language

Compiler
Technology

Performance Portable
Code Generator

OpenMPVHDL

C library...

Multicore CPUs

FPGAs

NN
Accelerator...

OpenCLGPUs...

Different starting point
Applications

Delite(top-down)

Image
Processing

NeuralNetworks

GraphAnalytics...

Data-Parallel
Intermediate Language

Performance Portable
Code Generator

Lift(bottom-up)

OpenMPVHDL

C library...

Multicore CPUs

FPGAs

NNAccelerator...

OpenCL

GPUs

What this talk is about
Applications
Image
ProcessingNeuralNetworks

GraphAnalytics...

Lift Data-Parallel
Intermediate Language

Lift rewrite-bases system+ hardware specific language

OpenMPVHDL

C library...

Multicore CPUs

FPGAs

NNAccelerator...

OpenCLGPUs

Rest of the talk►Part I: Lift data parallel language►

Part II: Optimisations as rewrite rules►

Part III: Code generation details

Part I
Lift: a Functional Data ParallelLanguage

Lift Intermediate Language
map(f) :⟼zip:⟼reduce(0,+):⟼split(n):⟼join:⟼iterate(f, n):⟼reorder(σ):⟼

Matrix Multiplication in Lift

Lift high-level matrix multiplication

Naive OpenCL version

Lift high-level matrix multiplication

Naive OpenCL version?

How to achieve high performance?

ARM Mali optimised version
Lift high-level matrix multiplication

Naive OpenCL version?

How to achieve high performance?

Part II
Encoding optimisations choices as
Rewrite Rules

Algorithmic Rewrite Rules• Provably correct rewrite rules• Express algorithmic and optimisations choices

Split-join rule:

Map fusion rule:
Reduce rules:...

mapmapsplit (m)mapjoinmapmapsplit (m)mapjoinmap/mforfor (int l = 0 ; l < m ; l++ ) {Map interchanged

Split-join ruleMap interchanged

A few rewrite steps later...

Tiled version

Low-Level OpenCL-specificExtensions

Lift OpenCL-Specific Primitives
map-global(f)map-workgroup(f)OpenCL thread hierarchy

map-local(f)map-sequential(f)Sequential implementations

reduce-sequential(z,f)
toLocal(f)Memory address spaces

toGlobal(f)vect-n(f)asScalarasVector-n

Vectorisation

OpenCL-specific rewrites(examples)
OpenCL thread hierarchy:

OpenCL memory hierarchy:

OpenCL vector types and operations:Let's start applying
OpenCL-specific rules

Vectorisation

Mapping parallelism to global threads

Accumulating in private memory

Result: high performanceNow just need to search the space...
100,000 runs later...

Performance Portability Achieved
Compiler input:

Search result:Desktop GPUs

Easily Extensible►

New rules can be added►

E.g. dot built-in
reduceSeq(z, add4) o mapSeq(mult4) o zip(x,y)=>
dot(x, y)Mali GPUEasily Extensible►

New rules can be added►

E.g. dot built-in
reduceSeq(z, add4) o mapSeq(mult4) o zip(x,y)=>
dot(x, y)Mali GPUPart III
Efficient Code Generation

Dot-product example

Dot-product example

after rewrite:Dot-product example

after rewrite:

copy to local
memoryDot-product example

after rewrite:

iterative
reduction in
local memoryDot-product example

after rewrite:
write back to
global mem.

How to generate efficient
OpenCL code?

→ Pattern based code generator
map-global (f,input)
for (uint i=get_global_id;i<n;i+= get_global_size) {output[i] = f(input[i]);parallel sum}

reduce-sequential (f,z,input)
T acc = z;
for (uint i=0; i<n; i++) {acc = f(acc, input[i]);}...

map-sequential (f,input)
for (uint i=0; i<n; i++) {output[i] = f(input[i]);}parallel sum

What about split, zip, join, … ?

►Need to avoid temporary results►

split, zip, … merely just change data layout•→ “lazy evaluation”

dot-product (a tiny bit of it):desired OpenCL output:
int wg_id = get_group_id(0) ;int l_id = get_local_id(0) ;...
acc = 0.0f;
for (int i = 0 ; i < 2 ; I += 1) {acc = acc +x [ 2 * l_id + 128 * wg_id + I ] *y [ 2 * l_id + 128 * wg_id + I ] ;}Construct a representation of theeffects of data layout functions:
Views
Consumes the views to generatecorrect array indicesView Construction

View Construction

View Construction

View Construction

View Construction

View Construction

View Construction

View Construction

View Consumption

View Consumption

View Consumption

View Consumption

View Consumption

View Consumption

View Consumption

View Consumption

View Consumption

dot-product (a tiny bit of it):produced OpenCL output:
int wg_id = get_group_id(0) ;int l_id = get_local_id(0) ;...
acc = 0.0f;
for (int i = 0 ; i < 2 ; I += 1) {acc = acc +x [ 2 * l_id + 128 * wg_id + I ] *y [ 2 * l_id + 128 * wg_id + I ] ;}

Array indices not so simple

Array indices not so simple

Can use well-known algebraic rules

Array indices not so simple

Can use well-known algebraic rules

and simplify array index expression

Putting it all together

rewritingcode generation

Code Generator vs Humanno simplification

arithmetic simplification

AMDNvidia

The Lift Approach: Summary►

Hardware-agnostic data-parallel functional intermediate language•►

Low-level functional language for each hardware type•►►hides hardware complexity, can be targeted by DSLs

OpenCL language (this talk), in the future: MPI, OpenMP

Rewrite rules based optimisation•

maps the hardware-agnostic language to the hardware-specific language•

extensible

Possible to achieve high performance through exploration•

opportunity for smarter technique (e.g. predictive modelling)

if you want to know more: www.lift-project.org
supported by: