Preparing for Extreme Heterogeneity
in High Performance Computing
Jeffrey S. Vetter
With many contributions from FTG Group and Colleagues

R-CCS International Symposium
Kobe
18 Feb 2020
ORNL is managed by UT-Battelle, LLC for the US Department of Energy

ORNL is managed by UT-Battelle
for the US Department of Energy

http://ft.ornl.gov

vetter@computer.org

Highlights• Recent trends in extreme-scale HPC paint an ambiguous future– Contemporary systems provide evidence that power constraints are driving architectures to change rapidly– Multiple architectural dimensions are being (dramatically) redesigned: Processors, node design, memory systems, I/O– Complexity is our main challenge• Applications and software systems are all reaching a state of crisis– Applications will not be functionally or performance portable across architectures– Programming and operating systems need major redesign to address these architectural changes– Procurements, acceptance testing, and operations of today’s new platforms depend on performance prediction andbenchmarking.

• We need portable programming models and performance prediction now more than ever!
– Heterogeneous processing• OpenACC->FGPAs• Intelligent runtime system (IRIS)• Clacc – OpenACC support in LLVM (not covered today)• OpenACC dialect of MLIR for Flang Fortran (not covered today)– Emerging memory hierarchies (NVM)• DRAGON – transparent NVM access from GPUs (not covered today)• NVL-C – user management of nonvolatile memory in C (not covered today)• Papyrus – parallel aggregate persistent storage (not covered today)

16• Performance prediction is critical for design and optimization (not covered today)

Time for a short poll…

28

History

Q: Think back 10 years. How
many of you would havepredicted that many of our
top HPC systems would be
GPU-based architectures?

30

Yes
No
Revisionists FutureAssume general purpose multicore CPU

Q: Think forward 10 years.
How many of you predictthat most of our top HPC
systems will have thefollowing architecturalfeatures?

GPU
FPGA/Reconfigurable processor
Neuromorphic processor
Deep learning processor
Quantum processor
RISC-V processor
Some new unknown processor
All/some of the above in one SoC

32

Implications
C, C++ XX, Fortran XX
Metaprogramming, etc (e.g., AMP, Kokkos, RAJA, SYCL)

Q: Now, imagine you are building
a new application with an
expected ~3M LOC and 20 teammembers over the next 10 years.
What on-node programming
model/system do you use?

CUDA, cu***, HIP, OpenCL
Directives: OpenMP XX, OpenACC XX
R, Python, Matlab, etc
A Domain Specific Language (e.g., Claw, PySL)
A Domain Specific Framework (e.g., PetSc)
Some new unknown programming approach
All/some of the above

34

The FTG Vision
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

Architectures
Multicore CPU

35

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

The FTG Vision | Applications
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

Architectures
Multicore CPU

36

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

https://exascaleproject.org/

ECP applications target national problems in 6 strategic areas
National security

Energy security

Economic security

Scientific discovery

Earth system

Health care

Stockpile
stewardship

Turbine wind plant
efficiency

Next-generationelectromagnetics
simulation of hostileenvironment andvirtual flight testing for
hypersonic re-entry
vehiclesHigh-efficiency,
low-emission
combustion engine
and gas turbine
designAdditive
manufacturingof qualifiable
metal parts

Find, predict,and control materials
and properties

Accurate regional
impact assessments
in Earth systemmodels

Accelerate
and translate
cancer researchMaterials design for
extremeenvironments ofnuclear fission
and fusion reactors
Design andcommercializationof Small Modular
Reactors
Subsurface use
for carbon capture,
petroleum extraction,
waste disposal
Scale-up of clean
fossil fuel combustion
Biofuel catalyst
design37

Reliable andefficient planningof the power grid
Seismic hazard
risk assessment
Urban planning

Cosmological probeof the standard modelof particle physics
Validate fundamental
laws of nature
Demystify origin ofchemical elements
Light source-enabled
analysis of protein
and molecular
structure and design
Whole-device modelof magneticallyconfined fusion
plasmasStress-resistant crop
analysis and catalytic
conversionof biomass-derived
alcohols
Metagenomics
for analysis ofbiogeochemical
cycles, climatechange,environmental
remediation

DARPA Domain Specific System on Chip Program is investigating Performance
Portability of Software Defined Radio
Xavier SoC #2

Xavier SoC #1

Video/Image
Files

UDP
Antenna
GR IEEE-802.11 Transmit (TX)

•

•
•
38

Signal processing: An opensource implementation of IEEE802.11 WIFI a/b/g with GR OOT
modules.
Input / Output file support via
Socket PDU (UDP server) blocks
Image/Video transcoding with
OpenCL/OpenCV

IEEE-802.11 Receive (RX)UDP

The FTG Vision | Architectures
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

Architectures
Multicore CPU

39

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

Contemporary devices are approaching fundamental limits

Economist, Mar 2016

Dennard scaling has already ended. Dennard observed that voltage andcurrent should be proportional to the linear dimensions of a transistor: 2xtransistor count implies 40% faster and 50% more efficient.

46

R.H. Dennard, F.H. Gaensslen, V.L. Rideout, E. Bassous, and A.R. LeBlanc, “Design of ion-implanted
MOSFET's with very small physical dimensions,” IEEE Journal of Solid-State Circuits, 9(5):256-68, 1974,

I.L. Markov, “Limits on fundamental limits to computation,” Nature, 512(7513):147-54,
2014, doi:10.1038/nature13570.

End of Moore’s Law : what’s your prediction ??

Economist, Mar 2016

47

“The number of people predicting the death of Moore’s Law doubles every two years.” – Peter Lee, Microsoft

48

Sixth Wave of Computing

6th wave

Transition
Periodhttp://www.kurzweilai.net/exponential-growth-of-computing

50

Predictions for Transition Period

51

Optimize Software andExpose New Hierarchical
Parallelism

Architectural
Specialization andIntegration• Redesign software toboost performance on
upcomingarchitectures• Exploit new levels ofparallelism andefficient data
movement• Use CMOS more
effectively for specificworkloads• Integrate components
to boost performanceand eliminate
inefficiencies• Workload specificmemory+storage
system design

Emerging Technologies• Investigate newcomputational
paradigms• Quantum• Neuromorphic• Advanced Digital• Emerging Memory
DevicesPredictions for Transition Period

52

Optimize Software andExpose New Hierarchical
Parallelism

Architectural
Specialization andIntegration• Redesign software toboost performance on
upcomingarchitectures• Exploit new levels ofparallelism andefficient data
movement• Use CMOS more
effectively for specificworkloads• Integrate components
to boost performanceand eliminate
inefficiencies• Workload specificmemory+storage
system design

Emerging Technologies• Investigate newcomputational
paradigms• Quantum• Neuromorphic• Advanced Digital• Emerging Memory
DevicesPredictions for Transition Period

53

Optimize Software andExpose New Hierarchical
Parallelism

Architectural
Specialization andIntegration• Redesign software toboost performance on
upcomingarchitectures• Exploit new levels ofparallelism andefficient data
movement• Use CMOS more
effectively for specificworkloads• Integrate components
to boost performanceand eliminate
inefficiencies• Workload specificmemory+storage
system design

Emerging Technologies• Investigate newcomputational
paradigms• Quantum• Neuromorphic• Advanced Digital• Emerging Memory
DevicesQuantum computing: Qubit design and fabrication
have made recent progress but still face challenges
Science 354, 1091 (2016) – 2 December54

http://nap.edu/25196

Fun Question: when was the field effect transistor patented?

https://www.edn.com/electronics-blogs/edn-moments/4422371/Lilienfeld-patents-field-effect-

58 transistor--October-8--1926

Predictions for Transition Period

59

Optimize Software andExpose New Hierarchical
Parallelism

Architectural
Specialization andIntegration• Redesign software toboost performance on
upcomingarchitectures• Exploit new levels ofparallelism andefficient data
movement• Use CMOS more
effectively for specificworkloads• Integrate components
to boost performanceand eliminate
inefficiencies• Workload specificmemory+storage
system design

Emerging Technologies• Investigate newcomputational
paradigms• Quantum• Neuromorphic• Advanced Digital• Emerging Memory
Devices

Pace of Architectural Specialization is Quickening• Industry, lacking Moore’s Law, will need to continue todifferentiate products (to stay in business)–

Use the same transistors differently to enhance performance• Architectural design will become extremely important,critical–Dark Silicon–

Address new parameters for benefits/curse of Moore’s Law• 50+ new companies focusing on hardware for Machine

http://www.theinquirer.net/inquirer/news/2477796/intels-nervanaai-platform-takes-aim-at-nvidias-gpu-techology

http://www.wired.com/2016/05/google-tpu-custom-chips/

Learning

D.E. Shaw, M.M. Deneroff, R.O. Dror et al., “Anton, a special-purpose machine for molecular dynamics
simulation,” Communications of the ACM, 51(7):91-7, 2008.

https://fossbytes.com/nvidia-volta-gddr6-2018/

HotChips 201860

HotChips 2018https://www.thebroadcastbridge.com/content/entry/1094/altera-announces-arria-10-2666mbps-ddr4-memory-fpga-interface

Xilinx ACAP

Analysis of Apple A-* SoCs

62

http://vlsiarch.eecs.harvard.edu/accelerators/die-photo-analysis

https://excl.ornl.gov/

Intel Stratix 10 FPGA

Experimental Computing Lab (ExCL) managed by the ORNL Future Technologies Group• Intel Stratix 10 FPGA and four banks of DDR4 external

memory–Board configuration: Nallatech 520 Network Acceleration Card

• Up to 10 TFLOPS of peak single precision performance• 25MBytes of L1 cache @ up to 94 TBytes/s peakbandwidth• 2X Core performance gains over Arria® 10• Quartus and OpenCL software

using FPGA

(Intel SDK v18.1) for• Provide researcher access to advanced FPGA/SOC

environment67

For more information or to apply for an account, visit https://excl.ornl.gov/

Mar 2019

https://excl.ornl.gov/

NVIDIA Jetson AGX Xavier SoC

Experimental Computing Lab (ExCL) managed by the ORNL Future Technologies Group• NVIDIA Jetson AGX Xavier:• High-performance system on a chip for autonomousmachines• Heterogeneous SoC contains:– Eight-core 64-bit ARMv8.2 CPU cluster (Carmel)– 1.4 CUDA TFLOPS (FP32) GPU with additional

inference optimizations (Volta)– 11.4 DL TOPS (INT8) Deep learning accelerator(NVDLA)– 1.7 CV TOPS (INT8) 7-slot VLIW dual-processor
Vision accelerator (PVA)– A set of multimedia accelerators (stereo, LDC,
optical flow)• Provides researchers access to advanced high-performance SOC environment

68

For more information or to apply for an account, visit https://excl.ornl.gov/

Mar 2019

Qualcomm 855 SoC (SM8510P) Snapdragon™
Experimental Computing Lab (ExCL) managed by the ORNL Future Technologies Grouphttps://excl.ornl.gov/

Qualcomm Development Board connected to (mcmurdo) HPZ820
7nm TSMC

Adreno 640

5G

Hexagon 690
Spectra 360
Kyro 485© Qualcomm Inc.
© Qualcomm Inc.

Connectivity (5G)

Kyro 485 (8-ARM Prime+BigLittle Cores)• Snapdragon X24 LTE (855 built-in) modem LTE Category 20

Prime
Core

• Snapdragon X50 5G (external) modem (for 5G devices)• Qualcomm Wi-Fi 6-ready mobile platform: (802.11ax-ready,802.11ac Wave 2, 802.11ay, 802.11ad)

Hexagon 690 (DSP + AI)• Quad threaded Scalar Core

• DSP + 4 Hexagon Vector Xccelerators• New Tensor Xccelerator for AI• Apps: AI, Voice Assistance, AV codecs
Adreno 640

69• Qualcomm 60 GHz Wi-Fi mobile platform: (802.11ay,
802.11ad)• Bluetooth Version: 5.0• Bluetooth Speed: 2 Mbps
• High accuracy location with dual-frequency GNSS.
Spectra 360 ISP• Vulkan, OpenCL, OpenGL ES 3.1• New dedicated Image Signal Processor (ISP)• Apps: HDR10+, HEVC, Dolby, etc• Dual 14-bit CV-ISPs; 48MP @ 30fps single camera• Enables 8k-360o VR video playback• Hardware CV for object detection, tracking, streo depth process• 20% faster compared to Adreno 630• 6DoF XR Body tracking, H265, 4K60 HDR video capture, etc.

• Connected Qualcomm board to HPZ820 through USB• Development Environment: Android SDK/NDK• Login to mcmurdo machine$ ssh –Y mcmurdo• Setup Android platform tools and development environment$ source /home/nqx/setup_android.source• Run Hello-world on ARM cores$ git clone https://code.ornl.gov/nqx/helloworld-android$ make compile push run• Run OpenCL example on GPU$ git clone https://code.ornl.gov/nqx/opencl-img-processing• Run Sobel edge detection$ make compile push run fetch• Login to Qualcomm development board shell$ adb shell$ cd /data/local/tmp

For more information or to apply for an account, visit https://excl.ornl.gov/

Created by Narasinga Rao Miniskar, Steve Moulton

Growing Open Source Hardware Movement Enables Rapid Chip Design

RISC-V Summit, 201870

DARPA ERI Programs Aiming for Agile (and Frequent) Chip Creation

A. Olofsson, 201871

Summary:
Transition Period will be Disruptive – Opportunities and Pitfalls Abound• New devices and architectures may not

be hidden in traditional levels of
abstraction• Examples– A new type of CNT transistor may becompletely hidden from higher levels– A new paradigm like quantum may require
new architectures, programming models, andalgorithmic approaches

Layer
Switch, 3D
Application
1
Algorithm
1
Language
1
API
1
Arch
1
ISA
1
Microarch
2
FU
2
Logic
3
Device3

NVM

Approximate
1
2
1
2
2
2
2
2
2
2
2
2
3
2
3
2
3
2
3
2

Adapted from IEEE Rebooting Computing Chart

72

Neuro

2
3
3
3
3
3
3
3
3
3

Quantum

3
3
3
3
3
3
3
3
3
3

Department of Energy (DOE) Roadmap to Exascale Systems
An impressive, productive lineup of accelerated node systems supporting DOE’s mission
Pre-Exascale Systems [Aggregate Linpack (Rmax) = 323 PF!]2012

Titan (9)ORNL

Cray/AMD/NVIDIA

2016

2018

IBM BG/Q

Cray/Intel KNL

AuroraANLIntel/Cray

NERSC-9
Perlmutter

Cray/IntelI/O
Xeon/KNL
Plateauing
Performance

103

AMD/Cray

ANL

LBNL

IBM BG/Q

ORNL

Deep Memory incl NVM

Cori (12)LLNL

2021-2023

IBM/NVIDIA

Theta (24)Sequoia (10)2020

Summit (1)Heterogeneous Cores ORNL

Mira (21)
ANLFirst U.S. Exascale Systems

Trinity (6)LANL/SNL

Cray/Intel Xeon/KNL

LBNL

Cray/AMD/NVIDIA

LLNL

Sierra (2)LLNL

IBM/NVIDIA

TBD

LANL/SNL
TBD

Jan 2018

Frontier Continues the Accelerated Node Design• Partnership between ORNL, Cray, and AMD• The Frontier system will be delivered in 2021
• Peak Performance greater than 1.5 EF• Composed of more than 100 Cray Shasta cabinets– Connected by Slingshot™ interconnect with adaptive routing, congestion control,and quality of service• Accelerated Node Architecture:– One purpose-built AMD EPYC™ processor– Four HPC and AI optimized Radeon Instinct™ GPU accelerators– Fully connected with high speed AMD Infinity Fabric links– Coherent memory across the node
– 100 GB/s injection bandwidth– Near-node NVM storage
118

Comparison of Titan, Summit, and Frontier Systems
System Specs

119

Titan

SummitFrontier

Peak

27 PF

200 PF~1.5 EF# cabinets

200

256

> 100

Node

1 AMD Opteron CPU
1 NVIDIA Kepler GPU

2 IBM POWER9™ CPUs
6 NVIDIA Volta GPUs

1 AMD EPYC CPU
4 AMD Radeon Instinct GPUs

On-nodeinterconnect

PCI Gen2
No coherence
across the node

NVIDIA NVLINK
Coherent memory
across the node

AMD Infinity Fabric
Coherent memory
across the node

SystemInterconnect

Cray Gemini network
6.4 GB/s

Mellanox Dual-port EDR IB network
25 GB/s

Cray four-port Slingshot network
100 GB/s

Topology

3D Torus

Non-blocking Fat Tree

Dragonfly

Storage

32 PB, 1 TB/s, Lustre
Filesystem

250 PB, 2.5 TB/s, IBM Spectrum
Scale™ with GPFS™

2-4x performance and capacityof Summit’s I/O subsystem.

On-node NVM

No

Yes

Yes

Power

9 MV

13 MV

29 MV

Complex architectures yield…

Complex
Programming
ModelsSystem: MPI, Legion, HPX, Charm++, etc
Low overhead

Resource contention

Locality120

Node: OpenMP, Pthreads, U-threads, etc
SIMD
NUMA, HBM

Cores: OpenACC, CUDA, OpenCL, OpenMP4, …
Memory use, coalescing

Data orchestration

Fine grained parallelism

Hardware features

During this Sixth Wave transition, Complexity is our major challenge!

DesignHow do we design future systems sothat they are better than current
systems on important applications?
• Simulation and modeling are more difficult• Entirely possible that the new system will be
slower than the old system!
• Expect ‘disaster’ procurements

121

Programmability
How do we design applications withsome level of performance portability?
• Software lasts much longer than transient
hardware platforms• Proper abstractions for flexibility andefficiency• Adapt or die

The FTG Vision | Programming Systems
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

Architectures
Multicore CPU

124

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

What more to say ?!?!? 

125

Directive-based Strategy with OpenARC: Open Accelerator
Research Compiler• Open-Sourced, High-Level Intermediate

Representation (HIR)-Based, Extensible
Compiler Framework.
–

Perform source-to-source translation from
OpenACC C to target accelerator models.
Support full features of OpenACC V1.0 ( + array
reductions and function calls)• Support both CUDA and OpenCL as target accelerator
models•– Provide common runtime APIs for various back-ends– Can be used as a research framework for various
study on directive-based accelerator computing.
Built on top of Cetus compiler framework, equipped withvarious advanced analysis/transformation passes andbuilt-in tuning tools.
• OpenARC’s IR provides an AST-like syntactic view of thesource program, easy to understand, access, andtransform the input program.
•

131

S. Lee and J.S. Vetter, “OpenARC: Open Accelerator Research Compiler for Directive-Based, Efficient Heterogeneous Computing,”
in ACM Symposium on High-Performance Parallel and Distributed Computing (HPDC). Vancouver: ACM, 2014

FPGAs| Approach• Design and implement an OpenACC-to-FPGA translation

framework, which is the first work to use a standard and portabledirective-based, high-level programming system for FPGAs.

• Propose FPGA-specific optimizations and novel pragma

extensions to improve performance.

• Evaluate the functional and performance portability of theframework across diverse architectures (Altera FPGA, NVIDIA
GPU, AMD GPU, and Intel Xeon Phi).

149

S. Lee, J. Kim, and J.S. Vetter, “OpenACC to FPGA: A Framework for Directive-based High-Performance Reconfigurable Computing,” Proc. IEEE
International Parallel & Distributed Processing Symposium (IPDPS), 2016, 10.1109/IPDPS.2016.28.

FPGA OpenCL Architecture
Host Processor

FPGA
PCIe

External DDR Memory

External DDR Memory

External Memory
Controller and PHY

External Memory
Controller and PHY

Global Memory Interconnect

Number of Replicated Compute Units
Pipeline
Depth
VectorWidth

KernelKernelPipeline
KernelPipeline
KernelPipeline
Pipeline

Local Memory
Interconnect

153

KernelKernelPipeline
KernelPipeline
KernelPipeline
Pipeline

Local Memory
Interconnect

KernelKernelPipeline
KernelPipeline
KernelPipeline
Pipeline

Local Memory
Interconnect

MemoryMemoryMemoryMemoryMemoryMemoryKernel-Pipelining Transformation Optimization• Kernel execution model in OpenACC– Device kernels can communicate witheach other only through the device
global memory.
– Synchronizations between kernels areat the granularity of a kernel
execution.
• Altera OpenCL channels– Allows passing data between kernels
and synchronizing kernels with highefficiency and low latency
156

Global Memory

Kernel 1

Kernel 2

Kernel communications throughglobal memory in OpenACC
Global Memory

Kernel 1 Channel

Kernel 2

Kernel communications with
Altera channels

Kernel-Pipelining Transformation Optimization (2)(a) Input OpenACC code#pragma acc data copyin (a) create (b) copyout (c){#pragma acc kernels loop gang worker present (a, b)for(i=0; i<N; i++) { b[i] = a[i]*a[i]; }#pragma acc kernels loop gang worker present (b, c)for(i=0; i<N; i++) {c[i] = b[i]; }}Global Memory

Kernel 1

Kernel 2(b) Altera OpenCL code with channels
channel float pipe_b;__kernel void kernel1(__global float* a) {int i = get_global_id(0);write_channel_altera(pipe_b, a[i]*a[i]);}__kernel void kernel2(__global float* c) {int i = get_global_id(0);c[i] = read_channel_altera(pipe_b);}
157

Global Memory

Kernel 1

ChannelKernel 2

Kernel-Pipelining Transformation Optimization (3)(a) Input OpenACC code#pragma acc data copyin (a) create (b) copyout (c){#pragma acc kernels loop gang worker present (a, b)for(i=0; i<N; i++) { b[i] = a[i]*a[i]; }#pragma acc kernels loop gang worker present (b, c)for(i=0; i<N; i++) {c[i] = b[i]; }}Global Memory

Kernel 1

Kernel 2

Kernel-pipelining
transformation

Valid under
specific conditions(c) Modified OpenACC code for kernel-pipelining#pragma acc data copyin (a) pipe (b) copyout (c){#pragma acc kernels loop gang worker pipeout (b) present (a)For(i=0; i<N; i++) { b[i] = a[i]*a[i]; }#pragma acc kernels loop gang worker pipein (b) present (c)For(i=0; i<N; i++) {c[i] = b[i];}}158

Global Memory

Kernel 1

ChannelKernel 2

159

Overall Performance of OpenARC FPGA Evaluation

FPGAs prefer applications with deep execution pipelines (e.g., FFT-1D and
FFT-2D), performing much higher than other accelerators.
For traditional HPC applications with abundant parallel floating-point operations,
it seems to be difficult for FPGAs to beat the performance of other accelerators,
even though FPGAs can be much more power-efficient.
• Tested FPGA does not contain dedicated, embedded floating-point
cores, while others have fully-optimized floating-point computation units.
Current and upcoming high-end FPGAs are equipped with hardened floatingpoint operators, whose performance will be comparable to other accelerators,
while remaining power-efficient.
175

The FTG Vision | Runtime and Operating Systems
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

Architectures
Multicore CPU

187

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

IRIS: Mapping Strategy for Heterogeneous Architectures and Native Programming
ModelsNVIDIA GPU

OpenMPARM CPU

OpenMP
OpenACCIntel OpenCL
Support moreprogramming
models.

SYCL

189

OpenCLHIP

IRIS Common Runtime API

CUDA

CPU/Xeon Phi
Intel FPGA
General
Accelerators
AMD GPU
IRIS offers a common API for diverse heterogeneous
devices and also allows intermixing of multiple
programming models (mix CUDA, OpenMP, OpenCL, etc.).

IRIS: An Intelligent Runtime System for Extremely HeterogeneousArchitectures• Provide programmers a unified programming

environment to write portable code acrossheterogeneous architectures (and preferred
programming systems)• Orchestrate diverse programming systems(OpenCL, CUDA, HIP, OpenMP for CPU) in a single
application– OpenCL• NVIDIA GPU, AMD GPU, ARM GPU, Qualcomm GPU, Intel
CPU, Intel Xeon Phi, Intel FPGA, Xilinx FPGA
– CUDA• NVIDIA GPU– HIP• AMD GPU
– OpenMP for CPU• Intel CPU, AMD CPU, PowerPC CPU, ARM CPU,
Qualcomm CPU

190

The IRIS Architecture• Platform Model– A single-node system equipped with host CPUs

and multiple compute devices (GPUs, FPGAs,
Xeon Phis, and multicore CPUs)• Memory Model– Host memory + shared device memory– All compute devices share the device memory• Execution Model– DAG-style task parallel execution across all

available compute devices• Programming Model– High-level OpenACC, OpenMP4, SYCL* (*planned)– Low-level C/Fortran/Python IRIS host-side
runtime API + OpenCL/CUDA/HIP/OpenMP
kernels (w/o compiler support)
191

Supported Architectures and Programming Systems by IRISExCL* Systems

OswaldSummit-nodeRadeonXavierSnapdragon

CPUIntel Xeon

IBM Power9Intel Xeon

ARMv8

Qualcomm
Kryo

Programming Systems

•
•

GPU
Programming Systems

192•

NVIDIA P100••

NVIDIA CUDA
NVIDIA OpenCL

IBM XL OpenMP••

NVIDIA V100•

NVIDIA CUDA

Intel OpenMPIntel OpenCL••

AMD Radeon VII••

Intel/Altera
Stratix 10

FPGA
Programming Systems

Intel OpenMPIntel OpenCLIntel OpenCL* ORNL Experimental Computing Laboratory (ExCL) https://excl.ornl.gov/

AMD HIPAMD OpenCL

GNU GOMP•

Qualcomm
Adreno 640

NVIDIA Volta•

NVIDIA CUDA

Android NDKOpenMP•

Qualcomm OpenCL

IRIS Booting on Various Platforms

193

Task Scheduling in IRIS
• A task– A scheduling unit
– Contains multiple in-order commands––•

Kernel launch command•

Memory copy command (device-to-host, host-to-device)

May have DAG-style dependencies with other tasks
Enqueued to the application task queue with a device
selection policy•

Available device selection policies– Specific Device (compute device #)– Device Type (CPU, GPU, FPGA, XeonPhi)– Profile-based– Locality-aware– Ontology-base– Performance models (Aspen)– Any, All, Random, 3rd-party users’ custom policies• The task scheduler dispatches the tasks in theapplication task queue to available compute devices–

194

Select the optimal target compute device according totask’s device selection policy

SAXPY Example on Xavier• Computation– S[] = A * X[] + Y[]• Two tasks– S[] = A * X[] on NVIDIA GPU (CUDA)– S[] += Y[] on ARM CPU (OpenMP)• S[] is shared between two tasks• Read-after-write (RAW), true dependency• Low-level Python IRIS host code +CUDA/OpenMP kernels– saxpy.py– kernel.cu– kernel.openmp.h

195

SAXPY: Python host code & CUDA kernel code
saxpy.py (1/2)#!/usr/bin/env python
import iris
import numpy as np
import sys
iris.init()

saxpy.py (2/2)
kernel0 = iris.kernel("saxpy0")
kernel0.setmem(0, mem_s, iris.iris_w)
kernel0.setint(1, A)
kernel0.setmem(2, mem_x, iris.iris_r)
off = [ 0 ]ndr = [ SIZE ]SIZE = 1024
A = 10.0

task0 = iris.task()
task0.h2d_full(mem_x, x)
task0.kernel(kernel0, 1, off, ndr)
task0.submit(iris.iris_gpu)

x = np.arange(SIZE, dtype=np.float32)
y = np.arange(SIZE, dtype=np.float32)s = np.arange(SIZE, dtype=np.float32)

kernel1 = iris.kernel("saxpy1")
kernel1.setmem(0, mem_s, iris.iris_rw)
kernel1.setmem(1, mem_y, iris.iris_r)

print 'X', x
print 'Y', y
mem_x = iris.mem(x.nbytes)
mem_y = iris.mem(y.nbytes)
mem_s = iris.mem(s.nbytes)

task1 = iris.task()
task1.h2d_full(mem_y, y)task1.kernel(kernel1, 1, off, ndr)
task1.d2h_full(mem_s, s)task1.submit(iris.iris_cpu)
print 'S =', A, '* X + Y', siris.finalize()196

kernel.cu (CUDA)
extern "C" __global__ void saxpy0(float* S, float
A, float* X) {int id = blockIdx.x * blockDim.x + threadIdx.x;
S[id] = A * X[id];}extern "C" __global__ void saxpy1(float* S,
float* Y) {int id = blockIdx.x * blockDim.x + threadIdx.x;
S[id] += Y[id];}SAXPY: Python host code & OpenMP kernel code
saxpy.py (1/2)#!/usr/bin/env python
import iris
import numpy as np
import sys
iris.init()

saxpy.py (2/2)
kernel0 = iris.kernel("saxpy0")
kernel0.setmem(0, mem_s, iris.iris_w)
kernel0.setint(1, A)
kernel0.setmem(2, mem_x, iris.iris_r)
off = [ 0 ]ndr = [ SIZE ]SIZE = 1024
A = 10.0

task0 = iris.task()
task0.h2d_full(mem_x, x)
task0.kernel(kernel0, 1, off, ndr)
task0.submit(iris.iris_gpu)

x = np.arange(SIZE, dtype=np.float32)
y = np.arange(SIZE, dtype=np.float32)s = np.arange(SIZE, dtype=np.float32)

kernel1 = iris.kernel("saxpy1")
kernel1.setmem(0, mem_s, iris.iris_rw)
kernel1.setmem(1, mem_y, iris.iris_r)

print 'X', x
print 'Y', y
mem_x = iris.mem(x.nbytes)
mem_y = iris.mem(y.nbytes)
mem_s = iris.mem(s.nbytes)

task1 = iris.task()
task1.h2d_full(mem_y, y)task1.kernel(kernel1, 1, off, ndr)
task1.d2h_full(mem_s, s)task1.submit(iris.iris_cpu)
print 'S =', A, '* X + Y', siris.finalize()197

kernel.openmp.h (OpenMP)#include <iris/iris_openmp.h>
static void saxpy0(float* S, float A, float* X,
IRIS_OPENMP_KERNEL_ARGS) {int id;#pragma omp parallel for shared(S, A, X)
private(id)IRIS_OPENMP_KERNEL_BEGIN
S[id] = A * X[id];IRIS_OPENMP_KERNEL_END}
static void saxpy1(float* S, float* Y,
IRIS_OPENMP_KERNEL_ARGS) {int id;#pragma omp parallel for shared(S, Y) private(id)IRIS_OPENMP_KERNEL_BEGIN
S[id] += Y[id];IRIS_OPENMP_KERNEL_END}

Memory Consistency Management
saxpy.py (1/2)#!/usr/bin/env python
import iris
import numpy as np
import sys
iris.init()

saxpy.py (2/2)
kernel0 = iris.kernel("saxpy0")
kernel0.setmem(0, mem_s, iris.iris_w)
kernel0.setint(1, A)
kernel0.setmem(2, mem_x, iris.iris_r)
off = [ 0 ]ndr = [ SIZE ]SIZE = 1024
A = 10.0

task0 = iris.task()
task0.h2d_full(mem_x, x)
task0.kernel(kernel0, 1, off, ndr)
task0.submit(iris.iris_gpu)

x = np.arange(SIZE, dtype=np.float32)
y = np.arange(SIZE, dtype=np.float32)s = np.arange(SIZE, dtype=np.float32)

kernel1 = iris.kernel("saxpy1")
kernel1.setmem(0, mem_s, iris.iris_rw)
kernel1.setmem(1, mem_y, iris.iris_r)

print 'X', x
print 'Y', y
mem_x = iris.mem(x.nbytes)
mem_y = iris.mem(y.nbytes)
mem_s = iris.mem(s.nbytes)

task1 = iris.task()
task1.h2d_full(mem_y, y)task1.kernel(kernel1, 1, off, ndr)
task1.d2h_full(mem_s, s)task1.submit(iris.iris_cpu)
print 'S =', A, '* X + Y', siris.finalize()

198

mem_s isshared between
GPU and CPULocality-aware Device Selection Policy
saxpy.py (1/2)#!/usr/bin/env python
import iris
import numpy as np
import sys
iris.init()

saxpy.py (2/2)
kernel0 = iris.kernel("saxpy0")
kernel0.setmem(0, mem_s, iris.iris_w)
kernel0.setint(1, A)
kernel0.setmem(2, mem_x, iris.iris_r)
off = [ 0 ]ndr = [ SIZE ]SIZE = 1024
A = 10.0

task0 = iris.task()
task0.h2d_full(mem_x, x)
task0.kernel(kernel0, 1, off, ndr)
task0.submit(iris.iris_gpu)

x = np.arange(SIZE, dtype=np.float32)
y = np.arange(SIZE, dtype=np.float32)s = np.arange(SIZE, dtype=np.float32)

kernel1 = iris.kernel("saxpy1")
kernel1.setmem(0, mem_s, iris.iris_rw)
kernel1.setmem(1, mem_y, iris.iris_r)

print 'X', x
print 'Y', y
mem_x = iris.mem(x.nbytes)
mem_y = iris.mem(y.nbytes)
mem_s = iris.mem(s.nbytes)

task1 = iris.task()
task1.h2d_full(mem_y, y)task1.kernel(kernel1, 1, off, ndr)
task1.d2h_full(mem_s, s)task1.submit(iris.iris_data)
print 'S =', A, '* X + Y', siris.finalize()

199

iris_data selects
the device that
requiresminimum datatransfer to
execute thetaskThe FTG Vision
Applications
Science andEngineering (e.g., CFD,
Materials, Fusion)

Streaming(e.g., SW Radio,
Experimental instrument)

Sensing(e.g., SAR, vision)

Deep learning(e.g., CNN)

Analytics(e.g., graphs)

Robotics(e.g., sense and react)

Compiler

Domain Specific
Languages

OpenARC

Just-in-time
Compilation

Metaprogramming

Scripting

Libraries

Autotuning

Synchronization

Load balancing

Runtime and Operating Systems
Task Schedulingand Mapping

Discovery

Data Orchestration

IO

IRIS

Architectures
Multicore CPU

201

GPU

FPGA

AI Accelerator

SoC

DSP

Deep Memory

Persistent
MemoryNeuromorphic

Performance
Productivity
Energy Efficiency

Programming Systems

Recap• Motivation: Recent trends in computing

paint an ambiguous future– Multiple architectural dimensions are being(dramatically) redesigned: Processors, nodedesign, memory systems, I/O– Complexity is our main challenge• Applications and software systems across

many areas are all reaching a state of crisis– Need a focus on performance portability• ORNL FTG investigating design andprogramming challenges for these trends– Performance modeling and ontologies– Performance portable compilation to manydifferent heterogeneous architectures/SoCs– Intelligent scheduling system to automate
discovery, device selection, and data movement– Targeting wide variety of existing and futurearchitectures (DSSoC and others)
204• Visit us– We host interns and other visitors yearround

• Faculty, grad, undergrad, high school,

industry• Jobs in FTG– Postdoctoral Research Associate in

Computer Science– Software Engineer– Computer Scientist– Visit https://jobs.ornl.gov• Contact me vetter@ornl.gov

Final Report on Workshop on Extreme Heterogeneity
1.

•

•

•

•

205

Maintaining and improving programmer productivity–

Flexible, expressive, programming models and languages–

Intelligent, domain-aware compilers and tools–

Composition of disparate software components

Managing resources intelligently–

Automated methods using introspection and machine learning–

Optimize for performance, energy efficiency, and availability

Modeling & predicting performance–

Evaluate impact of potential system designs and application mappings–

Model-automated optimization of applications

Enabling reproducible science despite non-determinism & asynchrony–

Methods for validation on non-deterministic architectures–

Detection and mitigation of pervasive faults and errors

Facilitating Data Management, Analytics, and Workflows–

Mapping of science workflows to heterogeneous hardware and software services–

Adapting workflows and services to meet facility-level objectives through learning approaches

https://orau.gov/exheterogeneity2018/https://doi.org/10.2172/1473756

Bonus Material

226