N-WAYS GPU BOOTCAMP
CUDA C/FORTRANCUDA C/FORTRAN
What to expect?
•

Basic introduction to GPU Architecture•

GPU Memory and Programming Model•

CUDA C/Fortran programming

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

GPU COMPUTING
Application CodeGPU

Rest of Sequential
CPU Code

Only Critical Functions
Parallelize using CUDA
Programming Model

+

CPU

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

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

CUDA KERNELS• Parallel portion of application: execute as a kernel• Entire GPU executes kernel, many threads• CUDA threads:• Lightweight• Fast switching• Tens of thousands execute simultaneously

CPU

Host

Executes functions

GPU

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

CUDA KERNELS: PARALLEL THREADS• A kernel is a function executed on theGPU• Array of threads, in parallel• All threads execute the same code, cantake different paths•

Each thread has an ID

float x = input[threadID];float y = func(x);output[threadID] = y;

• Select input/output data
• Control decisions

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

CUDA KERNELS: SUBDIVIDE INTO BLOCKS

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

CUDA KERNELS: SUBDIVIDE INTO BLOCKS•Threads are grouped into blocks

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

CUDA KERNELS: SUBDIVIDE INTO BLOCKS•Threads are grouped into blocks•Blocks are grouped into a grid9This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA KERNELS: SUBDIVIDE INTO BLOCKS•Threads are grouped into blocks•Blocks are grouped into a grid
•A kernel is executed as a grid of blocks of threads
10
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA KERNELS: SUBDIVIDE INTO BLOCKSGPU•Threads are grouped into blocks•Blocks are grouped into a grid
•A kernel is executed as a grid of blocks of threads
11
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

KERNEL EXECUTIONCUDA thread

CUDA core• Each thread is executed by a core

• Each block is executed by one SMand does not migrate
CUDA thread block

CUDA Streaming
Multiprocessor…

CUDA-capable GPU
CUDA kernel grid

...

…
…
…

• Several concurrent blocks canreside on one SM depending on
the blocks’ memory requirements
and the SM’s memory resources• Each kernel is executed on one
device• Multiple kernels can execute on a
device at one time

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

COMMUNICATION WITHIN A BLOCK•Threads may need to cooperate•Memory accesses•Share results•Cooperate using shared memory•Accessible by all threads within a block•Restriction to “within a block” permits scalability•Fast communication between N threads is not feasible when N large

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

TRANSPARENT SCALABILITY
1

2

3

4

5

6

7

8

9

10

11

12

9

10

7

8

5

6

3

4

1

2

11

12

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

TRANSPARENT SCALABILITY

1

2

3

4

5

6

7

8

9

10

9

10

11

12

1

2

3

4

11

12

5

6

7

8

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

TRANSPARENT SCALABILITY –
1

1

2

3

2

4

3

5

4

6

5

7

6

8

7

9

8

9

10

10

11

12

11

12

...
Idle
Idle Idle

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

CUDA PROGRAMMING MODEL - SUMMARY•A kernel executes as a grid ofthread blocks

DeviceHost

Kernel 1

0

1

2

3

0,0

0,1

0,2

0,3

1D

•A block is a batch of threads•Communicate through shared memory•Each block has a block ID

Kernel 2

2D
1,0

1,1

1,2

1,3

•Each thread has a thread ID
17
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA ARCHITECTURE
MEMORY MODEL
18
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

GPU ARCHITECTURE
Two Main components
Global memory
Analogous to RAM in a CPU server

HOST I/F
Giga Thread

Control units, registers, execution pipelines, caches

L2
DRAM I/F

Each SM has its own:

DRAM I/F

SMs perform the actual computations

DRAM I/F

DRAM I/F

Streaming Multiprocessors (SMs)

DRAM I/F

48GB with bandwidth currently up to 1 TB/s

DRAM I/F

Accessible by both GPU and CPU

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

MEMORY HIERARCHY•Thread•Registers

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

MEMORY HIERARCHY•Thread•Registers•Thread•Local memory

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

MEMORY HIERARCHY•Thread•Registers•Thread•Local memory•Block of threads•Shared memory

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

MEMORY HIERARCHY•Thread•Registers•Thread•Local memory•Block of threads•Shared memory•All blocks•Global memory
23
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA-C/FortranINTRODUCTION
24
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

WHAT IS CUDA?
•CUDA Architecture•Expose general-purpose GPU computing as first-class capability•Retain traditional DirectX/OpenGL graphics performance•CUDA C/Fortran•Based on industry-standard C/Fortran•A handful of language extensions to allow heterogeneous programs•Straightforward APIs to manage devices, memory, etc.

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

CUDA C/FORTRAN: THE BASICS▪ Terminology▪ Host – The CPU and its memory (host memory)▪ Device – The GPU and its memory (device memory)
Host

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

HELLO, WORLD!

int main( void ) {printf( "Hello, World!\n" );return 0;}

program main
implicit none
print *, “Hello World”
end program main

• This basic program is just standard C/Fortran that runs on the host
• NVIDIA’s compiler (nvcc/nvfortran) will not complain about CUDA programs with no device code
• At its simplest, CUDA C/Fortran is just C/Fortran!

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

HELLO, WORLD! WITH DEVICE CODE__global__ void kernel( void ) {}int main( void ) {kernel<<<1,1>>>();
printf( "Hello, World!\n" );return 0;}

module printgpu
contains
attributes(global) subroutine print_form_gpu()
implicit noneend subroutine print_form_gpu
end module printgpu
program testPrint
use printgpu
use cudaforimplicit none
call print_form_gpu<<<1, 1>>>()
end program testPrint• Two notable additions to the original “Hello, World!”
28
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

NVIDIA HPC SDK• Comprehensive suite of compilers, libraries, and tools used to GPU accelerate HPC
modeling and simulation application• The NVIDIA HPC SDK includes the new NVIDIA HPC compiler supporting CUDA C and Fortran•The command to compile CUDA C/C++ code is ‘nvcc’•

The command to compile CUDA Fortran code is ‘nvfortran’nvcc main.cu
nvfortran -cuda main.f90
29
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

HELLO, WORLD! WITH DEVICE CODE__global__ void kernel( void ) {}attributes(global) subroutine print_form_gpu()
end subroutine print_form_gpu•Keyword __global__ in C and attribute global in Fortran indicates that a function•Runs on the device•Called from host code

•nvcc/nvfortran splits source file into host and device components•NVIDIA’s compiler handles device functions like kernel()•Standard host compiler handles host functions like main()•gcc•Microsoft Visual C
30
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

HELLO, WORLD! WITH DEVICE CODE
int main( void ) {kernel<<< 1, 1 >>>();
printf( "Hello, World!\n" );return 0;}

program testPrint
use printgpu
use cudaforimplicit none
call print_form_gpu<<<1, 1>>>()
end program testPrint•Triple angle brackets mark a call from host code to device code
•Sometimes called a “kernel launch”•We’ll discuss the parameters inside the angle brackets later•This is all that’s required to execute a function on the GPU!
31
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

A MORE COMPLEX EXAMPLE•A simple kernel to add two integers:•As before, __global__ is a CUDA keyword meaning•add() will execute on the device … so a, b, and c must point to device memory•How do we allocate memory on the GPU?

•add() will be called from the host__global__ void add( int *a, int *b, int *c ) {*c = *a + *b;}

attributes(global) subroutine add( a, b, c)
int, device :: a(1), b(1), c(1)c(1) = a(1) + b(1)end subroutine vecAdd_kernel
32This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CPU + GPU

CPUPhysical Diagram▪

CPU memory is larger, GPU memory has more bandwidth▪

CPU and GPU memory are usually separate, connected by
an I/O bus (traditionally PCIe)▪

Any data transferred between the CPU and GPU will behandled by the I/O Bus▪
▪

The I/O Bus is relatively slow compared to memory
bandwidth
The GPU cannot perform computation until the data iswithin its memory

GPU

$

$

$

$

$

$

$

$

$

$

$

$

Shared Cache

$

$

$

$

$

$

$

$

Shared CacheHigh
Capacity
MemoryIO BusHigh Bandwidth
Memory33
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PROCESSING FLOW – STEP 1

PCI Bus

1. Copy input data from CPU memory to GPU
memory34
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PROCESSING FLOW – STEP 2

PCI Bus

1. Copy input data from CPU memory to GPU
memory2. Load GPU program and execute,caching data on chip for performance

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

PROCESSING FLOW – STEP 3

PCI Bus

1. Copy input data from CPU memory to GPU
memory2. Load GPU program and execute,caching data on chip for performance
3. Copy results from GPU memory to CPU
memory4. Unified Memory changes the nature of flow
• Some of the basics remains same

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

MEMORY MANAGEMENT•Host and device memory are distinct entities•Basic CUDA API for dealing with explicity device memory managment•cudaMalloc(), cudaFree(), cudaMemcpy()•Similar to their C equivalents, malloc(), free(), memcpy()•Similar to their Fortran equivalents, allocate(), deallocate•Array should be defined as allocatable•CUDA API for using Unified memory is• C API: cudaMallocManaged(), cudaFree()• Fortran: Declare variable with managed, allocatable attribute•real, managed, allocatable, dimension(:,:) :: A, B, C
* For this session we will be making use of Unified Memory
40
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

A MORE COMPLEX EXAMPLE: MAIN()
int main( void ) {int *a, *b, *c;
int size = sizeof( int );of a, b, c

program mainuse cudafor
real, managed, allocatable, dimension(:,:) ::a, b, ccudaMallocManaged( &b, size );allocate(a(n))
allocate(b(n))
allocate(c(n))

cudaMallocManaged( &c, size );call add<<<1, 1>>>(n, a, b, c)

add<<< 1, 1 >>>( a, b, c );deallocate(d_a)
deallocate(d_b)
deallocate(d_c)

cudaMallocManaged( &a, size );cudaFree( a ); cudaFree( b );end program main

cudaFree( c );return 0;}41This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PARALLEL PROGRAMMING IN CUDA
•But wait…GPU computing is about massive parallelism•So how do we run code in parallel on the device?
•Solution lies in the parameters between the triple angle brackets:

add<<< 1, 1 >>>( a, b, c );add<<< N, 1 >>>( a, b, c );•Instead of executing add() once, add() executed N times in parallel
42
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PARALLEL PROGRAMMING IN CUDA
•With add() running in parallel…let’s do vector addition•Terminology: Each parallel invocation of add() referred to as a block•Kernel can refer to its block’s index with the variable•C: blockIdx.x•Fortran: blockIdx%x•Each block adds a value from a[] and b[], storing the result in c[]:•By using blockIdx.x to index arrays, each block handles different indices43
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

PARALLEL PROGRAMMING IN CUDA__global__ void add( int *a, int *b, int *c){c[blockIdx.x] = a[blockIdx.x] +b[blockIdx.x];}

attributes(global) subroutine add(n, a, b, c)
integer, value :: nreal(8), device :: a(n), b(n), c(n)c(blockidx%x) = a(blockidx%x)+b(blockidx%x)}

end subroutine add

Block 0Block 1c[0] = a[0] + b[0];c[1] = a[1] + b[1];Block 2Block 3c[2] = a[2] + b[2];c[3] = a[3] + b[3];This is what runs in parallel on the device

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

THREADS•Terminology: A block can be split into parallel threads•Let’s change vector addition to use parallel threads instead of parallel blocks:•We use threadIdx.x instead of blockIdx.x in add()__global__ void add( int *a, int *b, int *c){c[threadIdx.x] = a[threadIdx.x] +b[threadIdx.x];}

attributes(global) subroutine add(n, a, b, c)
integer, value :: nreal(8), device :: a(n), b(n), c(n)c(threadidx%x) = a(threadidx%x)+b(threadidx%x)
end subroutine add

add<<<N,1>>>( a, b, c )add<<<1,N>>>( a, b, c );45
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

USING THREADS AND BLOCKS• We’ve seen parallel vector addition using•Many blocks with 1 thread apiece•1 block with many threads• Let’s adapt vector addition to use lots of both blocks and threads• After using threads and blocks together, we’ll talk about why threads• First let’s discuss data indexing…

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

INDEXING ARRAYS WITH THREADS AND BLOCKS•No longer as simple as just using threadIdx.x or blockIdx.x as indices•To index array with 1 thread per entry (using 8 threads/block)threadIdx.x

threadIdx.x

threadIdx.x

threadIdx.x

0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7
blockIdx.x = 0

blockIdx.x = 1

blockIdx.x = 2

blockIdx.x = 3•If we have M threads/block, a unique array index for each entry given by
int index = threadIdx.x + blockIdx.x * M;
int index =

x

+

y

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

INDEXING ARRAYS: EXAMPLE•In this example, the red entry would have an index of 21:
0

1

2

3

4

5

6

7

8

9 10 11 12 13 14 15 16 17 18 19 20 21

M = 8 threads/block

th

re

I
ad

dx

.x

=

5

blockIdx.x = 2

int index = threadIdx.x + blockIdx.x * M;
=

5

+

2

* 8;

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

ADDITION WITH THREADS AND BLOCKS• The blockDim.x is a built-in variable for threads per block:
int index= threadIdx.x + blockIdx.x * blockDim.x;• So what changes in main() when we use both blocks and threads?
attributes(global) subroutine add(n, a, b, c)__global__ void add( int *a, int *b, int *c ){int index = threadIdx.x + blockIdx.x* blockDim.x;}

c[index] = a[index] + b[index];

integer, value :: nreal(8), device :: a(n), b(n), c(n)integer :: id
id = (blockidx%x-1)*blockdim%x +threadidx%xc(id) = a(id) + b(id)end subroutine add

Note we used (blockidx%x-1) for Fortran as Fortran numbering starts from 1

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

PARALLEL ADDITION (BLOCKS/THREADS): MAIN()#define N(2048*2048)#define THREADS_PER_BLOCK 512

add<<< N/THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( a, b, c );blockSize = dim3(512,1,1)! Number of thread blocks in grid
gridSize = dim3(ceiling(real(n)/real(blockSize%x)) ,1,1)
call add<<<gridSize, blockSize>>>(n, d_a, d_b, d_c)50This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

CUDA SPEEDUP
Speed-up
4500.00X
4000.47X

4000.00X
3500.00X

Speed-Up

3000.00X
2500.00X
2000.00X
1500.00X
1000.00X
500.00X
0.00X

1.00X
SERIALNVIDIA TESLA V100
HPC SDK 20.11, NVIDIA Tesla V100, DGX1
51
This material is released by NVIDIA Corporation under the Creative Commons Attribution 4.0 International (CC BY 4.0)

REFERENCES

https://developer.nvidia.com/hpc-sdk

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

THANK YOU