Improving NAMD Performance on Volta GPUs David Hardy - Research - - PowerPoint PPT Presentation

Improving NAMD Performance

• n Volta GPUs

David Hardy - Research Programmer, University of Illinois at Urbana-Champaign Ke Li - HPC Developer Technology Engineer, NVIDIA John Stone - Senior Research Programmer, University of Illinois at Urbana-Champaign

Journey of a Legacy HPC Application

• NAMD has been developed for more than 20 years
• Parallel scaling of large systems on supercomputers
• Blue Waters (NCSA), Titan (ORNL), Stampede (TACC), Summit (ORNL)
• First full-featured molecular dynamics code to adopt CUDA
• Stone, et al. J Comput Chem, 28:2618-2640, 2007
• Why were certain design choices made?
• What lessons have we learned?
• Where does development need to go in the Age of Volta?

NAMD & VMD: Computational Microscope

NAMD - molecular dynamics simulation VMD - visualization, system preparation and analysis Enable researchers to investigate systems described at the atomic scale Ribosome Neuron Virus Capsid

1990 1994 1998 2002 2006 2010 104 105 106 107 108 2014 Lysozyme ApoA1 ATP Synthase STMV Ribosome HIV capsid Number of atoms 1986

Simulated System Sizes Have Increased Exponentially

Parallelism in Molecular Dynamics Limited to Each Timestep

Computational workflow of MD: Initialize coordinates forces, coordinates Update coordinates Force calculation Occasional output of reduced quantities (energy, temperature, pressure) Occasional output of coordinates (trajectory snapshot)

about 99% of computational work about 1% of computational work

Work Dominated by Nonbonded Forces

90% — Non-bonded forces, short-range cutoff 5% — Long-range electrostatics, gridded (e.g. PME) 2% — Bonded forces (bonds, angles, etc.) 2% — Correction for excluded interactions 1% — Integration, constraints, thermostat, barostat

force calculation update coordinates Apply GPU acceleration first to the most expensive part

Parallelize by Spatial Decomposition of Atoms

Data parallelism is common to many codes

NAMD Also Decomposes Compute Objects

• Spatially decompose data and

communication

• Separate but related work

decomposition

• “Compute objects” create

much greater amount of parallelization, facilitating iterative, measurement-based load balancing system, all from use of Charm++

Kale et al., J. Comp. Phys. 151:283-312, 1999

Phillips et al., SC2002

Offload to GPU

Objects are assigned to processors and queued as data arrives

Overlap Calculations, Offload Nonbonded Forces

Early Nonbonded Forces Kernel Used All Memory Systems

• Start with most expensive calculation: direct nonbonded interactions.
• Decompose work into pairs of patches, identical to NAMD structure.
• GPU hardware assigns patch-pairs to multiprocessors dynamically.

16kB Shared Memory

Patch A Coordinates & Parameters

32kB Registers

Patch B Coords, Params, & Forces

Texture Unit

Force Table Interpolation

Constants

Exclusions 8kB cache 8kB cache 32-way SIMD Multiprocessor 32-256 multiplexed threads

768 MB Main Memory, no cache, 300+ cycle latency

Force computation on single multiprocessor (GeForce 8800 GTX has 16)

NAMD Performance Improved Using Early GPUs

• Full NAMD, not test harness
• Useful performance boost

– 8x speedup for nonbonded – 5x speedup overall w/o PME – 3.5x speedup overall w/ PME – GPU = quad-core CPU

• Plans for better performance

– Overlap GPU and CPU work. – Tune or port remaining work.

• PME, bonded, integration, etc.

ApoA1 Performance

faster seconds per step 0.75 1.5 2.25 3 CPU GPU Other PME Nonbond

2.67 GHz Core 2 Quad Extreme + GeForce 8800 GTX

Reduce Communication Latency by Separating Work Units

Remote Force Local Force GPU CPU Other Nodes/Processes Local Remote x f f f f Local x x Update One Timestep x

Phillips et al., SC2008

Early GPU Fits Into Parallel NAMD as Coprocessor

• Offload most expensive calculation: non-bonded forces
• Fits into existing parallelization
• Extends existing code without modifying core data structures
• Requires work aggregation and kernel scheduling considerations to
• ptimize remote communication
• GPU is treated as a coprocessor

Biomedical Technology Research Center for Macromolecular Modeling and Bioinformatics Beckman Institute, University of Illinois at Urbana-Champaign - www.ks.uiuc.edu

GTC 2017

NAMD Scales Well on Kepler Based Computers

(2fs timestep)

0.25 0.5 1 2 4 8 16 32 64 256 512 1024 2048 4096 8192 16384 Number of Nodes 21M atoms 224M atoms Performance (ns per day) Blue Waters XK7 (GTC16) Titan XK7 (GTC16) Edison XC30 (SC14) Blue Waters XE6 (SC14)

14

Kepler based

Large Rate Difference Between Pascal and CPU

• Balance between GPU and CPU capability keeps shifting towards GPU
• NVIDIA plots show only through Pascal — Volta widens the performance gap!
• Difference made worse by multiple GPUs per CPU (e.g. DGX, Summit)
• Past efforts to balance work between GPU and CPU are now CPU bound

20x FLOP rate difference between GPU and CPU Requires full use of CPU cores and vectorization!

Case Study: Multilevel Summation CUDA Kernels in VMD

Electrostatic field of chromatophore model from multilevel summation method: computed with 3 GPUs (G80) in ~90 seconds, 46x faster than single CPU socket

• Effort to balance computational work between GPUs and CPU
• GPU gets only straightforward data parallel algorithms, regularized work units
• CPU gets algorithms less well suited to GPU, plus “overflow” work

Hardy, et al. Journal of Parallel Computing, 35:164-177, 2009

Partial model: ~10M atoms

Le, et al. PLoS Comput Bio, 6:e1000939, 2010

Balancing Work Between GPU and CPU

Time-Averaged Electrostatics Analysis on 
 NCSA Blue Waters

Preliminary performance for VMD time-averaged electrostatics w/ Multilevel Summation Method on the NCSA Blue Waters Early Science System

NCSA Blue Waters Node Type Seconds per trajectory frame for

• ne compute node

Cray XE6 Compute Node: 32 CPU cores (2xAMD 6200 CPUs) 9.33 Cray XK6 GPU-accelerated Compute Node: 16 CPU cores + NVIDIA X2090 (Fermi) GPU 2.25 Speedup for GPU XK6 nodes vs. CPU XE6 nodes XK6 nodes are 4.15x faster overall Tests on XK7 nodes indicate MSM is CPU-bound with the Kepler K20X GPU. Performance not much faster (yet) than Fermi X2090 Need to move spatial hashing, prolongation, interpolation onto the GPU… In progress…. XK7 nodes 4.3x faster overall

speedup ok CPU-bound

Multilevel Summation on the GPU

Computational steps CPU (s) w/ GPU (s) Speedup Short-range cutoff 480.07 14.87 32.3 Long-range anterpolation 0.18 restriction 0.16 lattice cutoff 49.47 1.36 36.4 prolongation 0.17 interpolation 3.47 Total 533.52 20.21 26.4

Performance profile for 0.5 Å map of potential for 1.5 M atoms. Hardware platform is Intel QX6700 CPU and NVIDIA GTX 280.

Accelerate short-range cutoff and lattice cutoff parts

Multilevel summation of electrostatic potentials using graphics processing

• units. D. Hardy, J. Stone, K. Schulten. J. Parallel Computing, 35:164-177, 2009.

Case Study: Multilevel Summation CUDA Kernels in VMD

• Successful (at the time) of exploiting task and data level parallelism
• Work was reasonably well balanced between CPUs and GPUs for Tesla and Fermi
• For Kepler and later, approach was CPU bound

Conclusions:

Balancing Work Between GPU and CPU

Reduce Latency, Offload All Force Computation

• Overlapped GPU communication and computation (2012)
• Offload atom-based work for PME (2013)
• Use higher order interpolation with coarser grid
• Reduce parallel FFT communication
• Faster nonbonded force kernels (2016)
• Offload entire PME using cuFFT (for single node use) (2016)
• Offload remaining force terms (2017)
• Includes: bonds, angles, dihedrals, impropers, crossterms, exclusions

Emphasis on improving communication latency Emphasis on using GPUs more effectively

Overlapped GPU Communication and Computation

• Allows incremental results from a single grid to be

processed on CPU before grid finishes on GPU

• Allows merging and prioritizing of remote and local work
• GPU side:
• Write results to host-mapped memory (also without streaming)
• __threadfence_system() and __syncthreads()
• Atomic increment for next output queue location
• Write result index to output queue
• CPU side:
• Poll end of output queue (int array) in host memory

Non-overlapped Kernel Communication

Integration unable to start until GPU kernel finishes

Overlapped Kernel Communication

GPU kernel communicates results while running; patches begin integration as soon as data arrives

7 S6623: Advances in NAMD GPU Performance

Non-bonded force computation in NAMD

• Two levels of spatial

sorting

– Simulation box is divided into patches – Within the patch, atoms are sorted spatially into groups of 32 using

• rthogonal recursive

bisection method

4 4 4

Faster

8 S6623: Advances in NAMD GPU Performance

Non-bonded force compute

• Compute = all pairwise

interactions between two patches

Compute 1 Compute 2 32 32 Compute 1 Patch 1 Patch 2 Patch 3 Patch 2

• For GPU,

compute is split into tiles of size 32x32 atoms

Faster

9 S6623: Advances in NAMD GPU Performance

Non-bonded force computation

Atoms in patch i Atoms in patch j

32 31 30 1 2 3

Fj Fi

• One warp per tiles
• Loop through 32x32 tile diagonally

– Avoids race condition when storing forces Fi and Fj

• Bitmask used for exclusion lookup

32 32 Warp 1 Warp 2 Warp 3 Warp 4

Faster

11 S6623: Advances in NAMD GPU Performance

Neighbor list construction on GPU

Bounding box neighbor list Compute forces – atom-based neighbor list Sort neighbor list

R R

Sort neighbor list

12 S6623: Advances in NAMD GPU Performance

Neighbor list sorting

• Tile lists executed on the same thread block should

have approximately the same work load

• Simple solution is to sort according to tile list length
• Also minimizes tail effects at the end of kernel

execution

Load imbalance! No load imbalance Thread block sort

Warp 1

Warp 2

Warp 3

Warp 1 Warp 2 Warp 3 Warp 4

Single-Node GPU Performance Competitive on Maxwell

New kernels by Antti-Pekka Hynninen, NVIDIA

Stone, Hynninen, et al., International Workshop on OpenPOWER for HPC (IWOPH'16), 2016

More Improvement from Offloading Bonded Forces

• GPU offloading for bonds, angles, dihedrals,

impropers, exclusions, and crossterms

• Computation in single precision
• Forces are accumulated in 24.40 fixed point
• Virials are accumulated in 34.30 fixed point
• Code path exists for double precision

accumulation on Pascal and newer GPUs

• Reduces CPU workload and hence improves

performance on GPU-heavy systems

DGX-1 Speedup

1 1.175 1.35 1.525 1.7

apoa1 f1atpase stmv

New kernels by Antti-Pekka Hynninen, NVIDIA

Supercomputers Increasing GPU to CPU Ratio

Blue Waters, Titan with Cray XK7 nodes 1 K20 / 16-core AMD Opteron Summit nodes 6 Volta / 42 cores IBM Power 9

Only 7 cores supporting each Volta!

Limited Scaling Even After Offloading All Forces

Results on NVIDIA DGX-1 (Intel Haswell using 28-cores with Volta V100 GPUs)

0.5 1 1.5 2 2.5 3 3.5 4 1 2 3 4 STMV 1 million atoms Performance (ns per day) Number of NVIDIA Voltas Offloading all forces Nonbonded forces only

NAMD 2.12 NAMD 2.13

CPU Integrator Calculation (1%) Causing Bottleneck

Nsight Systems profiling of NAMD running STMV (1M atoms) on 1 Volta & 28 CPU cores GPU is not being kept busy

Too much communication! [nonbonded] [bonded] [PME]

Too much CPU work: 2200 patches across 28 cores Patches running sequentially within each core

CPU integrator work is mostly data parallel, but…

• Uses double precision for positions, velocities, forces
• Data layout is array of structures (AOS), not well-suited to vectorization
• Each NAMD “patch” runs integrator in separate user-level thread to

make source code more accessible

• Benefit from vectorization is reduced, loop over 200–600 atoms in each patch
• Too many exceptional cases handled within same code path
• E.g. fixed atoms, pseudo-atom particles (Drude and lone pair)
• Test conditionals for simulation options and rare events (e.g. trajectory output)

every timestep

CPU integrator work is mostly data parallel, but…

• Additional communication is required
• Send reductions for kinetic energies and virial
• Receive broadcast periodic cell rescaling when simulating constant pressure
• A few core algorithms have sequential parts, with reduced parallelism
• Not suitable for CPU vectorization
• Rigid bond constraints — successive relaxation over each hydrogen group
• Random number generator — Gaussian numbers using Box–Muller transform

with rejection

• Reductions calculated over “hydrogen groups” — irregular data access patterns

Strategies for Overcoming Bottleneck

• Data structures for CPU vectorization
• Convert atom data storage from AOS (array of structures) form into vector friendly SOA

(structure of arrays) form

• Algorithms for CPU vectorization
• Replace non-vectorizing RNG code with vectorized version
• Replace rigid bond constraints sequential algorithm with one capable of fine-grained

parallelism (maybe LINCS or Matrix-SHAKE)

• Offload integrator to GPU
• Main challenge is aggregating patch data
• Use vectorized algorithms, adapt curand for Gaussian random numbers

Goal: Developing GPU-based NAMD

• CPU primarily manages GPU kernel launching
• CPU prepares and aggregates data structures for GPU, handles Charm++ communication
• Reduces overhead of host-to-device memory transfers
• Data lives on the GPU, clusters of patches
• Communicate edge patches for force calculation and atom migration
• Design new data structures capable of GPU or CPU vectorization
• Major refactor of code to reveal more data parallelism
• Kernel-based design with consistent interfaces across GPU and CPU
• Need fallback kernels for CPU (e.g. continue to support Blue Waters and Titan)

Some Conclusions

For HPC apps in general and NAMD in particular

• Balance between CPU and GPU computational capability continues to shift in favor
• f the GPU
• The 1% workload from 10 years ago is now 60% or more in wall clock time
• Any past code that has attempted to load balance work between CPU and GPU is

today likely to be CPU bound

• Best utilization of GPU might require keeping all data on the GPU
• Motivates turning a previously CPU-based code using GPU as an accelerator into a GPU-

based code using CPU as a kernel management and communication coprocessor

• Volta + CUDA 9.x + CUDA Toolkit Libraries provide enough general purpose

support to allow moving an application entirely to the GPU

Related Talks

• S8747 - ORNL Summit: Petascale Molecular Dynamics Simulations on the

Summit POWER9/Volta Supercomputer - James Phillips

• S8709 - Accelerating Molecular Modeling Tasks on Desktop and Pre-Exascale

Supercomputers - John Stone

• S8718 - Optimizing HPC Simulation and Visualization Codes Using the NVIDIA

Nsight Systems - Robert Knight, Daniel Horowitz, John Stone

• S8665 - VMD: Biomolecular Visualization from Atoms to Cells Using Ray Tracing,

Rasterization, and VR - John Stone

• SE150572 - OpenACC User Group Meeting - Sunita Chandrasekaran, Guido

Juckeland, John Stone, Randy Allen

Acknowledgments

• Special thanks to:
• NVIDIA CUDA Team
• NVIDIA NSight Systems Team
• James Phillips, NCSA, UIUC
• In memoriam Antti-Pekka Hynninen, NVIDIA
• Theoretical Biophysics Research Group, Beckman Institute, UIUC
• Parallel Programming Laboratory, Dept of Computer Science, UIUC
• Grants:
• NIH P41-GM104601 Center for Macromolecular Modeling and Bioinformatics
• Summit Center for Accelerated Application Readiness, OLCF, ORNL