2D-Heat

2D-Heat equation < CUDA implementation>

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

Definition Mathematische Definition (2D): Beschreibt den Temperaturverlauf der - Fläche über der Zeit . Anwendungen in: Finanz Mathematik Statistik Physik, Chemie

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

Diskretisierung Diskretisierung des Gebietes Finite Differenzen Methode (FDM)

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

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 zu lösenden heat equation

heat.cu void heat() { // Kernel Konfiguration dim3 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); // Pointeraustausch tmp = d_src; d_src = d_dst; d_dst = tmp; } }

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 Kernels intx = blockIdx.x * blockDim.x + threadIdx.x; inty = blockIdx.y * blockDim.y + threadIdx.y; inti = y * width + x; if(x < width - 2 && y < height - 2) { float sum = 0.0; // oben sum += 0.25 * d_src[i + 1]; // links sum += 0.25 * d_src[i + width]; // rechts sum += 0.25 * d_src[i + width + 2]; // unten sum += 0.25 * d_src[i + 2 * width + 1]; // Mitte schreiben d_dst[i + width + 1] = sum; } }

heat_kernel.cu (shared) __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(); // Rechnen float 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 }

Optimierung • Speicherzugriffe zw. Host u. Device • Durch umbiegen der Zeiger minimiert • Shared 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 maximiert • Beste Kernel Konfiguration für unser Problem gefunden

Messungen Global Memory

Messungen Shared Memory

Global vs. Shared

GPU vs. CPU Single Core (faui07g) : Intel Pentium 4 → 2.4 GHz Dual Core (faui08) : AMD Opteron → 2.2 GHz Quad Core (faui06i): Intel Q6600 → 2.4 GHz

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

Demo • Sascha …

→ → ? →

2D-Heat

2D-Heat

Presentation Transcript

2D coordinates

2D Drawings

2D Shapes.

2D Sketches

2D shapes

Animação 2D

HEAT HEAT HEAT

2d Division

2D hexagonal

2D shape

2D Viewing

2D Transformations

Soal-2d

Labo 2d

WinIso 2D

Clipping 2D

2D Transformations

2D Transformations

Jasełka 2D

2D Viewing

2D Photography

2D ANIMATION