1 2D-Heat equation. 2 Überblick Definition Diskretisierung Iterationsverfahren Mapping auf Cuda...

1

2D-Heat

< CUDA implementatio

n>

equation

2

Überblick

Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo

3

Definition Mathematische Definition (2D):

Beschreibt den Temperaturverlauf der - Fläche über der Zeit .

Anwendungen in: Finanz Mathematik Statistik Physik, Chemie

4

Überblick


5

Diskretisierung Diskretisierung des Gebietes

Finite Differenzen Methode (FDM)

6

Überblick


7

Iterationsverfahren Einsetzen in die DGL ergibt finiten Differenzenstern

Filtern des „Bildes“ mit diesem Filterkernel ergibt einen Zeitschritt

Berechnungsverfahren: Jacobiverfahren gut parallelisierbar Gauß-Seidel Verfahren schlecht parallelisierbar

8

Überblick


9

Mapping auf Cuda Durch Jacobiverfahren können viele Pixel unabhängig

voneinander berechnet werden

Kernel (Thread) berechnet ein gefiltertes Pixel Benötigt 4 umliegende Werte und sich selbst

Grid berechnet eine Iteration (einen Filterdurchlauf)

Bild wird in Blöcke zerlegt

Anzahl der Iterationen (Filteraufrufe)bestimmt die simulierte Zeit der zulösenden heat equation

10

Überblick


11

heat.cu

void heat() {// Kernel Konfigurationdim3 block(BLOCK_SIZE_X, BLOCK_SIZE_Y, 1);dim3 grid(width / block.x, height / block.y, 1);

// Filteraufruf for(i = 0; i < iterations; i++) {

filter<<<grid, block>>>(d_src, d_dst, width, height);filter<<<grid, block>>>(d_src, d_dst, width, height);

// Pointeraustauschtmp = d_src;d_src = d_dst;d_dst = tmp;

}}

12

heat_kernel.cu (global)

__global__ void filter_global(float* d_src, float* d_dst, int width, int height) {

// Referenz ist linke obere Ecke des Filter Kernelsint x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;int i = y * width + x;

if(x < width - 2 && y < height - 2) {float sum = 0.0;// obensum += 0.25 * d_src[i + 1];// linkssum += 0.25 * d_src[i + width];// rechtssum += 0.25 * d_src[i + width + 2];// untensum += 0.25 * d_src[i + 2 * width + 1];

// Mitte schreibend_dst[i + width + 1] = sum;

}}

13

__shared__ float img_block[BLOCK_SIZE_X + 2][BLOCK_SIZE_Y + 2];

img_block[tx + 1][ty + 1] = d_src[i]; // Mitte lesen

__syncthreads();

if(x > 0 && x < width - 1 && y > 0 && y < height - 1) {//Bildrand nicht if(tx == 0) { verändern

img_block[tx][ty + 1] = d_src[i - 1]; // links} else if(tx == blockDim.x - 1) {

img_block[tx + 2][ty + 1] = d_src[i + 1]; // rechts}if(ty == 0) {

img_block[tx + 1][ty] = d_src[i - width]; // oben} else if(ty == blockDim.y - 1) {

img_block[tx + 1][ty + 2] = d_src[i + width]; // unten}

__syncthreads(); // Rechnenfloat sum = 0.0;sum += 0.25 * img_block[tx][ty + 1];sum += 0.25 * img_block[tx + 2][ty + 1];sum += 0.25 * img_block[tx + 1][ty];sum += 0.25 * img_block[tx + 1][ty + 2];d_dst[i] = sum; // Zurückschreiben

}

heat_kernel.cu (shared)

14

Überblick


15

Optimierung

Speicherzugriffe zw. Host u. DeviceSpeicherzugriffe zw. Host u. Device Durch umbiegen der Zeiger minimiert

Shared memory verwendetShared memory verwendet Coalscaled Speicherzugiff

- 32 * 4 Byte = 128 Byte werden pro Warp linear gelesen- Startadresse des source-images ist ganzzahliges Vielfaches von 128

Rechnen auf Shared Memory schneller

Occupancy maximiertOccupancy maximiert Beste Kernel Konfiguration für unser Problem gefunden

16

Überblick


17

Messungen Global Memory

8X – 8Y

Bildgröße Zeit [ms]

64 x 6464 x 64 0.0139128 x 128128 x 128 0.0532256 x 256256 x 256 0.2200512 x 512512 x 512 0.9000

1024 x 10241024 x 1024 3.60002048 x 20482048 x 2048 14.4000

16X – 8Y


64 x 6464 x 64 0.0129128 x 128128 x 128 0.0478256 x 256256 x 256 0.1897512 x 512512 x 512 0.7576

1024 x 10241024 x 1024 3.05002048 x 20482048 x 2048 12.2200

16X – 16Y


64 x 6464 x 64 0.0130128 x 128128 x 128 0.0500256 x 256256 x 256 0.1900512 x 512512 x 512 0.7560

1024 x 10241024 x 1024 3.03002048 x 20482048 x 2048 12.2700

18

Messungen Shared Memory

8X – 8Y


64 x 6464 x 64 0.0083128 x 128128 x 128 0.0280256 x 256256 x 256 0.1160512 x 512512 x 512 0.4540

1024 x 10241024 x 1024 1.80002048 x 20482048 x 2048 7.3000

16X – 8Y


64 x 6464 x 64 0.0044128 x 128128 x 128 0.0096256 x 256256 x 256 0.0260512 x 512512 x 512 0.0940

1024 x 10241024 x 1024 0.36902048 x 20482048 x 2048 1.4700

16X – 16Y


64 x 6464 x 64 0.0044128 x 128128 x 128 0.0100256 x 256256 x 256 0.0270512 x 512512 x 512 0.0980

1024 x 10241024 x 1024 0.38002048 x 20482048 x 2048 1.5540

19

Global vs. Shared

Global Memory (16X – 8Y)


64 x 6464 x 64 0.0129128 x 128128 x 128 0.0478256 x 256256 x 256 0.1897512 x 512512 x 512 0.7576

1024 x 10241024 x 1024 3.05002048 x 20482048 x 2048 12.2200

Shared Memory (16X – 8Y)


64 x 6464 x 64 0.0044128 x 128128 x 128 0.0096256 x 256256 x 256 0.0260512 x 512512 x 512 0.0940

1024 x 10241024 x 1024 0.36902048 x 20482048 x 2048 1.4700

20

GPU vs. CPU

Device GFLOPS

Single CoreSingle Core ~ 0.084

Dual CoreDual Core ~ 0.384

Quad CoreQuad Core ~ 0.225

GPUGPU ~ 22.78

Single Core (faui07g)Single Core (faui07g) : Intel Pentium 4 → 2.4 GHzDual Core (faui08)Dual Core (faui08) : AMD Opteron → 2.2 GHzQuad Core (faui06i):Quad Core (faui06i): Intel Q6600 → 2.4 GHz

Speedup

59x

101x

273x

0

50

100

150

200

250

300

GPU vs. Dual Core GPU vs. Quad Core GPU vs. Single Core

21

Überblick


22

Probleme

Bildgrößen nur als Vielfaches von Blockgröße möglich

Zeiten messen Timer in cuda hat inkonsistente Zeiten geliefert Profiler liefert bessere Ergebnisse

Teilweise Branching und Bankkonflikte nicht vermeidbar

23

Überblick


24

Demo Sascha …

25

0,0000

2,0000

4,0000

6,0000

8,0000

10,0000

12,0000

14,0000

64x64 128x128 256x256 512x512 1024x1024 2048x2048

Zeit

[ms]

→

→→?

1 2D-Heat equation. 2 Überblick Definition Diskretisierung Iterationsverfahren Mapping auf Cuda...

Documents

Transcript of 1 2D-Heat equation. 2 Überblick Definition Diskretisierung Iterationsverfahren Mapping auf Cuda...