Post on 05-Apr-2015
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
5
Diskretisierung Diskretisierung des Gebietes
Finite Differenzen Methode (FDM)
6
Überblick
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Bildgröße Zeit [ms]
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
Bildgröße Zeit [ms]
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
Bildgröße Zeit [ms]
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
Bildgröße Zeit [ms]
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
Bildgröße Zeit [ms]
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)
Bildgröße Zeit [ms]
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)
Bildgröße Zeit [ms]
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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
Definition Diskretisierung Iterationsverfahren Mapping auf Cuda Implementierung Optimierung Performance Vergleiche Probleme Demo
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]
→
→→?