INTRODUCTION TO
GPU COMPUTING
N-WAYS GPU BOOTCAMP
12This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

GPU PROGRAMMING(FOUNDATIONS)

3This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

FUNDAMENTALLY DIFFERENTGPU Accelerator
Optimized for
Parallel Tasks

4This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

SILICON BUDGETLess

ALUMoreMoreControl

Less

MoreCacheLess

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

5CPU IS A LATENCY REDUCING ARCHITECTURE
Strengths
GPUCPU
Accelerator

CPUOptimized for
Serial Tasks

Optimized for• Very large
main memory
ParallelTasks• Very fast clock speeds• Latency optimized via large caches• Small number of threads can runvery quickly
CPU Weaknesses• Relatively low memory bandwidth• Cache misses very costly• Low performance/watt6

GPU IS ALL ABOUT HIDING LATENCYGPU Strengths• High bandwidthCPU
main memory

Optimizedfor• Significantly
more computeresources Serial Tasks• Latency tolerant via parallelism• High throughput• High performance/watt

GPU Accelerator
Optimized forParallel Tasks

GPU Weaknesses• Relatively low memory capacity• Low per-thread performance

7

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

HETEROGENEOUS PROGRAMMING

8

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

EXECUTION FLOW – H2D

9This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

EXECUTION FLOW – KERNEL LAUNCH

10This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

EXECUTION FLOW – D2H

11This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

MANY WAYS TO PROGRAM A GPU(PROGRAMMING MODELS)

17This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

HOW GPU ACCELERATION WORKS
Application CodeCompute-Intensive Functions

GPURest of Sequential
CPU Code

5% of Code+CPU

18

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

GPU PROGRAMMING IN 2020 AND BEYOND
Math Libraries | Standard Languages | Directives | CUDA__global__
void saxpy(int n, float a,float *x, float *y) {int i = blockIdx.x*blockDim.x +threadIdx.x;
if (i < n) y[i] += a*x[i];}#pragma acc data copy(x,y){
std::transform(par, x, x+n, y, y,[=](float x, float y) {return y + a*x;});...

do concurrent (i = 1:n)y(i) = y(i) + a*x(i)
enddo

...

GPU Accelerated
C++ and Fortranint main(void) {cudaMallocManaged(&x, ...);
cudaMallocManaged(&y, ...);...
saxpy<<<(N+255)/256,256>>>(...,x, y)
cudaDeviceSynchronize();...
}

std::transform(par, x, x+n, y, y,[=](float x, float y) {return y + a*x;});}Incremental Performance
Optimization with Directives

Maximize GPU Performance with
CUDA C++/FortranGPU Accelerated Math Libraries
19

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

HPC DEVELOPERS NEED A ON-RAMP

20

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

BOOTCAMP MINI-APP
28

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

APPLICATION
Molecular Simulation

RDF
The radial distribution function (RDF) denoted inequations by g(r) defines the probability offinding a particle at a distance r from anothertagged particle.

29

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

RDF

Pseudo Code - Cfor (int frame=0;frame<nconf;frame++){for(int id1=0;id1<numatm;id1++){for(int id2=0;id2<numatm;id2++){
dx=d_x[]-d_x[];dy=d_y[]-d_y[];dz=d_z[]-d_z[];r=sqrtf(dx*dx+dy*dy+dz*dz);
if (r<cut) {ig2=(int)(r/del);d_g2[ig2] = d_g2[ig2] +1 ;}Across Frames

Find DistanceReduction}}}30

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

RDF

Pseudo Code - Fortrando iconf=1,nframes
if (mod(iconf,1).eq.0) print*,iconf
do i=1,natoms
do j=1,natoms
dx=x(iconf,i)-x(iconf,j)dy=y(iconf,i)-y(iconf,j)dz=z(iconf,i)-z(iconf,j)
r=dsqrt(dx**2+dy**2+dz**2)
if(r<cut)theng(ind)=g(ind)+1.0d0
endif
enddo
enddo
enddo

Across Frames

Find Distance

Reduction

31

This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

32