CPU vs. GPU

Floating-point performance

• En av grunnene er at de har svært god ytelse i forhold til CPUer.

• GPUene er først å fremst raskere en CPuene, men vi kan også se at
• ytelsen til GPUer øker raskere enn ytelsen til CPUer.
• De har en større stigende kurve...
• Det er viktig å huske at dette er den teoretiske beste ytelsen, og at det er skal veldig mye til for å få til en så god ytelsen…

Memory bandwidth

• En annen fordel er at GPUer har høy minnebåndbredde...
• Ofte er applikasjoner bundet av minne, så en høy minnebåndbredde er viktig for å oppnå høy ytelse for mange problemer....

Design

• En CPU har noen få kjerner (typisk to, fire eller åtte kjerner) som er veldig komplekse og raske.
• En GPU har mange (hundrevis) enkle og trege kjerner.

• Hver CPU kjerne har egen avansert kontroll logikk.
• Hos en GPU så deler grupper av kjerner kontrol logikk.

• En CPU skjuler minne ventetid ved bruk av “cacheing” og “prefetching”.
• En GPU skjuler minne ventetid ved å bytte mellom tråder som kjører.
• Hvis en tråd venter på minne aksess, så bytter GPUen til annen tråd, som kan gjøre beregninger istedenfor.
• Det å bytte mellom tråder går veldig fort (man har veldig rask kontekst svitsjing).

Transistors

• CPU: Uses large fraction of the chip area for control logic and cache
• GPU: Uses most of the chip area for data processing units

• Oppsumert, så er GPUen utformet slik at flere transistorer er viet til å gjøre beregninger enn til data caching og flytkontroll.
• De grønne boksene er prosesserings enheter, mens de gule og oransje boksene er data caching og flytkontroll.
• En CPU er designet for å kjøre noen få tråder veldig raskt.
• En GPU er designet for å kjøre tusenvis av tregere tråder samtidig.

Latency vs. Throughput

What is better? What do you need?

• En CPU kan sammenlignes med to sportsbiler.
• En GPU kan sammenlignes med mange scootere.

• En CPU er optimalisert for lav ventetid, altså designet for å minimere kjøretiden til en enkel tråd.
• En GPU er optimalisert for høy gjennomstrømning, altså designet for å få prossesert ferdig flest mulig tråder innenfor en gitt tid.

• Beregn en jobb så fort som mulig -> CPU, kan sammenlignes med f.eks levere en pakke så snart som mulig.
• Beregn mange jobber innenfor en gitt tid -> GPU, kan sammenlignes med f.eks levere mange pakker innenfor en gitt tid.

• Det er mange applikasjoner hvor det å optimalisere for høy gjennomstrømning er mest viktig.
• I f.eks bilde prosessering er man mer opptatt av antall piksler per sekund, enn ventetiden av en hvilken som helst piksel.
• Man er villig til å la tiden det tar å prosessere en piksel ta dobbel så lang tid, så lenge vi får prosessert flere pikseler på innenfor en git tid.
• For dette vil en GPU være gunstig.

• Hvor lang tid tar det å bli ferdig med en oppgave
• Time (Seconds)

• Antall oppgaver ferdig i en gitt tid
• Stuff / Time (Jobs/Hours)
• Image processing - Pixels / Seconds

• Da skal vi begynne å se på hvordan vi kan programmere NVIDIA sine GPUer...
• Det blir en del teori først nå, før vi skal ta for oss ett kode eksempel (hvor vi skal legge sammen to vektorer)...

What is CUDA?

• Compute Unified Device Architecture
• Nvidia introduced CUDA in 2006
• Specially designed for GPGPU
• Only supported by Nvidia graphics cards
• Write code in C/C++, Java, Python, Fortran, Perl...
• You do not need parallel programming experience
• Requires no knowledge of graphics APIs
• Access to native instructions and memory

Get started - CUDA Toolkit

• Driver
• Compiler nvcc
• Development, profiling and debugging tools
• Various libraries
• Programming guides, and API reference
• Example codes

• Før man kan starte å programmere med CUDA, så må man laste ned “CUDA Toolkit”...
• Inneholder en del biblioteker som kan vurdere å bruke....
• Eksempel koden er veldig fint å se på....

Get started

• 1) På venstre her har vi det å bruke biblioteker:
  • cuFFT, - Fast Fourier Transforms Library
  • cuBLAS - Complete BLAS bibliotek
  2) I midten har vi OpenACC som er svært lik OpenMP programmering, for de som har gjort det.
  • Man legger til PRAGMAer (direktiver) i koden, som forteller kompilatoren at den skal parallisere den kommenterte delen av koden for deg.
  • F.eks hvis man skriver #pragma acc parallel foran en for-løkker, skal kompliatoren automatisk paralllisere for-løkken for deg...
  3) Men i dag skal vi skal bruke programmeringsspråk, som er den mest kraftfulle og fleksible måten. Vi vil bruke C som språk...

CUDA C Programming Guide

The advent of multicore CPUs and many-core GPUs means that mainstream processor chips are now parallel systems. Furthermore, their parallelism continues to scale with Moore’s law. The challenge is to develop application software that transparently scales its parallelism to leverage the increasing number of processor cores...

At its core are three key abstractions – a hierarchy of thread groups, shared memories, and barrier synchronization...

...data parallelism...

CUDA terminology

Heterogeneous programming

• Det vi snakker om i dag er GPU programmering, også kalt heterogen programmering.
• Som er koordinering av to eller flere forskjellige prosessorer, av forskjellige arkitektur typer, for å utføre en felles oppgave.
• Man kombinerer det beste fra to verdener: CPU + GPU
• CPU og GPU er en kraftfull kombinasjon fordi CPUer er optimalisert for seriell prosessering, mens GPU består er optimalisert for parallell prosessering.
• Parallelle deler av et program kjøres på GPUen som “kernels”.

Kernels

C functions, that when called, are executed N times in parallel by N different threads on the device. As opposed to only once, like a regular C functions.

• det er en C-funksjon, som, når den blir kalt, blir utført N ganger i parallell av N forskjellige tråder, i motsetning til en vanlig C-funksjon, som bare blir utført en gang.
• ...så alt som skjer inne funksjonen, blir utført N antall ganger...
• __global__ definerer at funksjonen er en kernel, og at funksjonen skal bli kjørt på GPUen...
• Antallet ganger “kernelen” blir utført (eller antall tråder), er spesifisert med en spesiell syntaks… <<<>>>...vi skal komme tilbake til det snart...

Thread hierarchy

• En “kernel” blir utført av tusenvis av tråder i parallell..
• ...i form av ett hierarki!

Thread hierarchy - Blocks

• En “kernel” blir utført av tusenvis av tråder i parallell..
• ...i form av ett hierarki!
• Som vi ser på bildet...
• Tråder er gruppert i blokker (1D, 2D, eller 3D)
• Hver tråd har en unik lokal ID i blokken sin...

Thread hierarchy - Grids

• Blokker er gruppert i et grid (1D, 2D, eller 3D).
• ….så vi har altså ett grid at blokker, hvor hver blokk inneholder mange tråder!
• Hver blokk har en unik ID i gridet...
• Alle blokker i et grid har samme dimensjon...

Thread hierarchy - Scalability

• Det at trådblokker er uavhengig, gjør at koden man skriver, kan skaleres med antall GPU kjerner…
• slik at en GPU med flere kjerner vil kjøre programmet på raskere enn en GPU med færre kjerner...

How to launch a kernel?

Execution configuration example

dim3 dimGrid(2, 2); // 4 blocks (2D)
dim3 dimBlock(2, 4); // 8 threads per block (2D)
myKernel<<<dimGrid, dimBlock>>>();

dim3 is a structure with three properties, x, y and z.

• SEE SLIDE
• ...mellom trippel vinkelbraketter
• Number of threads = 4 * 8 = 32 threads
• Det er bare ett eksempel på hvordan man kan starte en kernel, vanligvis har man mange flere tråder….

How to determine unique thread IDs?

With this execution configuration

dim3 dimGrid(3); // 3 blocks in 1D
dim3 dimBlock (5); // 5 threads per block in 1D

How are kernels executed?

• The GPU core is the stream processor (SP)
  • Able to run a single sequential thread
• SPs are grouped into streaming multiprocessors (SMs)
  • Can execute hundreds of threads concurrently
  • SMs is basically a SIMD processor
• There are multiple SMs per GPU

• For å forklarer hvordan GPUer utfører “kernels”, må vi først se på hvordan GPUer fra NVIDIA er designet ….
• Hver stream prosessor er i stand til å kjøre en sekvensiell tråd...
• Stream processorer er gruppert til streaming multiprocessorer...
• En streaming multiprosessor kan kjøre et stort antall tråder samtidig, hvor hver tråd kjører det samme program.
• En stream multiprosessor kan ses på som en SIMD-prosessor.

NVIDIA's consumer graphics cards

How are kernels executed - Grid

• Ett grid er håndert av en enkel GPU.
• Det er GPUen som er ansvarlig for allokering av blokker til streaming multiprocessors med tilgjengelig kapasitet… Kan også nevne at dette blir gjort automatisk av GPUen...

How are kernels executed - Block

• En blokk blir håndert av bare en enkel streaming multiprocessor.

How are kernels executed - Thread

• En stream processor kan håndtere en eller flere tråder i en blokk.

Synchronization

• Point in the program where threads stop and wait
• When all threads have reached the barrier, they can proceed

Inside kernel

__syncthreads() // Synchronize threads within a block

• Funksjons deklarasjoner, for å angi om en funksjon skal eksikveres på CPuen eller på GPUen.
• Variabel deklarasjoner, for å angi minne plassering på GPUen.
• Du har funksjon for å synkronisere tråder innenfor en trådblokk.

Memory hierarchy

• Tråder kan få tilgang til data fra flere forskjellige minneområder på GPUen under eksekvering, illustrert i figuren...
• Hver tråd har ett privat lokalt minneområde.
• Hver trådblokk har ett delt minneområde som er synlig for alle tråder i blokken.
• Minnområde har samme levetid som trådblokken.
• Alle tråder har tilgang til ett stort globalt minneområde.
• Ytelse (husk å holde dataene i den raskeste minne som er overhode mulig):
• Lokalt -> Delt -> Global

Memory management

Inside kernel

float variable; // Local memory (registers)
__shared__ float variable; // Shared memory
__device__ float variable; // Global memory

Memory management

• Allocate GPU memory:
  • cudaMalloc(......)
• Copy data to/from GPU:
  • cudaMemcpy(......, cudaMemcpyHostToDevice)
  • cudaMemcpy(......, cudaMemcpyDeviceToHost)
• Free GPU memory:
  • cudaFree(......)

CUDA C Programming Guide...

At its core are three key abstractions – a hierarchy of thread groups, shared memories, and barrier synchronization...

• Vi har ett tråd hiarki….
• Vi har ett minne hiarki…
• Vi har synkroniserings muligheter…

CUDA example - Vector addition

Typical program execution

• Allocate memory and initialize data on CPU
• Allocate memory on GPU
• Transfer data from CPU to GPU
  • A slow operation, aim to minimize this!
• Lunch kernel
• Transfer results back from GPU to CPU
• Free CPU and GPU memory

The problem to solve

C = A + B

Vector addition on CPU

for(int i=0; i<N-1; i++) {
    C[i] = A[i] + B[i];
}

• Only one thread of execution
• No explicit parallelism

• Det er bare en tråd som kjører
• Det er ingen eksplisitt parallellitet

Allocate memory on CPU

int main() {
    // Size of vectors
    int n = 50000;

    // Size, in bytes, of each vector
    size_t bytes = n * sizeof(float);

    // Allocate memory for the host vectors
    float *h_A = (float *) malloc(bytes);
    float *h_B = (float *) malloc(bytes);
    float *h_C = (float *) malloc(bytes);

• For GPU eksempelet skal vi først skal se på main funksjonen, for så å se på den interessante biten, nemlig kernelen.....
• Data on the host (CPU) starts with h_
• Data on the device (GPU) starts with d_

Initialize vectors with some values

for(int i=0; i<n; ++i) {
      h_A[i] = 1;
      h_B[i] = 3;
    }

Allocate memory on GPU

// Allocate memory for the device vectors
    float *d_A = NULL;
    cudaMalloc((void **) &d_A, bytes);
    float *d_B = NULL;
    cudaMalloc((void **) &d_B, bytes);
    float *d_C = NULL;
    cudaMalloc((void **) &d_C, bytes);

Transfer data from CPU to GPU

// Copy the host vectors to the device
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

Lunch kernel

// Execution configuration
    int threadsPerBlock = 256;
    int blocksPerGrid =(n + threadsPerBlock - 1) / threadsPerBlock;

    // Launch kernel
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B,
    d_C, n);

    printf("Kernel launch with %d blocks of %d threads\n",
    blocksPerGrid, threadsPerBlock);

Transfer results back from GPU to CPU

// Copy the result back to the host result vector
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

Show the result

// Sum up host result vector and print result divided by n,
    // this should equal 4
    float sum = 0;
    for(int i=0; i<n; i++) {
        sum += h_C[i];
    }
    sum = sum / n;
    printf("Final result: %f\n", sum);

Free CPU and GPU memory

// Free device global memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

The kernel

The for-loop is removed!

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void vectorAdd(float *d_A, float *d_B, float *d_C, int n) {
    // Calculate the thread ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // Make sure we do not go out of bounds
    if (i < n) {
        d_C[i] = d_A[i] + d_B[i];
    }
}

• De tre vektorene A, B og C blir tatt inn som parametere, og er lagret i det globalet minnet.
• N spesifiserer antall elementer i vektorene, og er lagret i det lokalet minnet til tråden.
• Hvis vi ser på GPU eksempelet, så ser vi at for-løkken har blitt borte.
• Istedenfor blir denne “kernelen” kjørt i parallell av mange tråder, hvor hver tråd operere på ett spesifikt element fra hver vektor.
• Vi bruker tråd IDen for å bestemme hvilket element en tråd skal operere på.
• Vi har minst like mange tråder som det er elementer i vektorene.

Visualized

• Oppsumert, så ser vi her hvordan GPU “kernelen” legger sammens de to vektorene A og B…
• Vi har like mange tråder som det er vektor elementer…
• Tråd indeksen blir brukt til å bestemme hvilket elementer en tråd skal legge sammen….
• Det viktigste å legge merke til i eksempelet er SIMT-tankegangen.

Example: Vector addition on GPU

$ Kernel launch with 196 blocks of 256 threads
$ Final result: 4.000000

Execution time CPU vs. GPU - Assume we do 64 additions

• CPU uses 2 ns for 1 addition. CPU execution time = number of additions * time it takes for 1 addition = 64 * 2 ns = 128 ns
• GPU uses 10 ns for 1 addition. GPU execution time = the time it takes for 1 addition (we assume we have enough resources to do the all additions in parallel) = 10 ns

• Hvis vi ser på ett veldig forenklet eksempel på hvordan kjøretiden mellom en CPU. vs GPU…
• I akkurat dette eksempelet så er kjøretiden til GPUen mindre, selv om den bruker lengre tid på en adisjon, i forhold til CPUen...

Final comments

Performance guidelines

• High arithmetic intensity: Math (maximize) / memory (minimize)
• High number of threads: Keep the GPU busy. Need enough threads to hide memory latency
• Avoid thread divergence (caused by if, switch, do, for, while statements): Different execution paths -> serialisation (SIMT)
• Large data sets: Plenty to operate on in parallel
• Minimize CPU-GPU memory transfers: Slow operation
• Use single precision if possible
• ...

• Aritmetisk intensitet - Sier noe om hvor mange regne operasjoner et program utfører i forhold til antall minne operasjoner.
• Det er ønskelig at denne er høy (for å skjule minne latens).

• Maksimere antall bergeninger per tråd
• Minimere antall minne aksesser
• Minimere tid brukt på minne aksess
• Bruk raskest mulig minne

Why use GPUs?

• Computing power. It is powerful
  • Number crunching: 1 GPU ~= 4 TFLOPS ~= Small cluster
• It is a cheap commodity hardware
  • FLOPS per NOK
• It is everywhere
  • Sold hundreds of millions programmable GPUs
• Power
  • FLOPS per Watt

NVIDIA - Tesla cards

• No graphic output! What?
• Aimed at scientific computing rather than gaming
• Tested and burned-in for long-running calculations

• ...kan ikke brukes som grafikkort...det finnes ingen grafikk utgang....
• Laget for generelle beregninger...
• BILDE FORKLARING: Her ser du et tesla kort som du kan putt inn i arbeidsstasjonen din, eller man kan kjøpe en server løsning…
• ...man kan selvsakt bruke vanlige NVIDIA grafikk kort for å gjøre CUDA programmering...

GPGPU - Areas

• Graphics
• Multimedia
• Ultrasound Imaging
• Molecular Dynamics
• Seismic Imaging
• Astrophysic
• Data Mining
• Finance
• Physics
• Chemistry
• ...

More Responsibility

• GPU programming puts more responsibility on the programmer than "regular" programming
• For example the programmer must decide and explicitly code:
  • How to partition the computation into a grid, blocks, and threads
  • Where (in what kind of memory) to put each piece of data for best performance

• GPU-programmering legger mer ansvar på programmereren enn “vanlig” programmering.
• En programmerer må bestemme og eksplisitt kode:
• Hvordan dele opp en beregning i ett grid, blokker og tråder….
• Hvor (i hva slags minne) man bør plassere data for best ytelse….
• osv...

The end!

• CPU + GPU = combination of flexibility and performance
• GPUs can make your simulations go faster. So that you can not slack off....
• Getting started with GPU programming is easy, being able to fully utilize GPU hardware is hard...
• CPU optimizations does not apply to GPU
• Memory movement is very expensive

References

• State of the art virtualized graphics
• GPU programming using CUDA
• CUDA Programming
• Introduction to GPU computing
• Tutorial on GPU computing - With an introduction to CUDA. Felipe A. Cruz
• CU - CSFP - 130320 - GPU Computing and CUDA
• Udacity - Intro to Parallel Programming