Hauptseminar MAP08 Random Heightmap on GPU Hannes Stadler, Sebastian Graf

1. Hauptseminar MAP08 Random Heightmap on GPU Hannes Stadler, Sebastian Graf HannesStadler@gmx.de, sebgr@gmx.net Betreuung: Matthias Hartl, Hritam Dutta, Frank Hannig Hardware-Software-Co-Design Universität Erlangen-Nürnberg

2. Gliederung • Was ist eine Heightmap? • Fault Algorithmus • Parallelisierung des Fault Algorithmus • Umsetzung in Cuda • Benchmarks • Probleme • Zusammenfassung

3. Was ist eine Heightmap? • dt. Höhenfeld • Zwei-dimensionales Skalarfeld • Beschreibung eines Höhenreliefs • Jedem Punkt ist ein Wert zugeordnet, der dessen Höhe angibt

4. Fault Algorithmus • Erzeuge ein ebenes 2-dimensionales Grid Algorithmus: • Wähle zwei zufällige Punkte im R² • Lege Gerade durch diese zwei Punkte • Erhöhe alle Punkte auf der eine Seite der Gerade, erniedrige die auf der anderen um einen konstanten Wert • Wiederhole diese Schritte für eine vorher festgelegt Anzahl von Iterationen

5. Fault Algorithmus • Pseudo-Code: foreach(Iteration) { CreateRandomLine(); foreach( RowOfImage) { foreach(PixelOfRow) { processNewValue(); } } }

6. Fault Algorithmus

7. Variationen des Fault Algorithmus • Multiplikation der Geraden mit Sinus/Cosinus um weiche Übergänge an den Kanten zu bekommen

8. Parallelisierung Fault Algorithmus • Algorithmus besteht aus drei for-Schleifen • Parallelisierung der Schleifen: • 1.for-Schleife durchläuft Anzahl der Iterationen  mehrere Iterationen parallel möglich, da unabhängig • 2.for-Schleife führt Berechung für jede Zeile im Bild aus Parallelisierbar, da Zeilen unabhängig • 3.for-Schleife arbeitet auf genau einer Zeile Pro Zeile eine Grenze ( Schnittpunkt mit der Geraden), Aufteilung in Teil der erhöht und der erniedrigt wird

9. Umsetzung in CUDA • Naiver Ansatz: • Laden des Grids in Global Memory • Threads arbeiten auf Daten im Global Memory • Probleme mit Nebenläufigkeit, Performance etc. • Optimierter Ansatz: • Aufteilung des Grids in Blöcke • Block in Shared Memory laden • Berechnung aller Iterationen für jeweiligen Block • Danach wieder zurück in Global Memory speichern • Weitere Optimierungen: • Coalesced Speicherzugriff der Threads • Zugriff auf Zufallzahlen über Constant Memory • Oder: Zufallszahlen auf der GPU erzeugen

10. Quellcode – Kernelaufruf int CreateHeightMap() { CUT_DEVICE_INIT(); dim3 threads(TPL,ZPB);// 16 x 16 dim3 grid(WIDTH/BLOCKWIDTH,HEIGHT/ZPB); // 1k x 1k -> 64 x 8, bei 2k x 2k -> 128 x 16 int rand[ITERATIONS*4]; for(i=0;i<ITERATIONS*4;i++){ // rand[ ] mit Zufallszahlen füllen } CUDA_SAFE_CALL(cudaMemcpyToSymbol(rand_d, rand ,ITERATIONS*4*sizeof(int),0) ); GLfloat* HeightMap_d; CUDA_SAFE_CALL(cudaMalloc((void**) &HeightMap_d, WIDTH*HEIGHT*sizeof(float))); splitpicture<<<grid, threads>>>(HeightMap_d); CUDA_SAFE_CALL(cudaMemcpy(HeightMap, HeightMap_d , WIDTH*HEIGHT*sizeof(float), cudaMemcpyDeviceToHost) ); }

11. Quellcode – Kernel extern __constant__ int rand_d[]; __global__ void splitpicture(GLfloat *HeightMap_d) { __shared__ float aRow[ZPB][SMB/ZPB]; // Init. mit default-Wert weggelassen for(int i = 0; i < ITERATIONS; i++){ // Variableninitialisierung, random-Werte, „Wendestelle“ bestimmen for(int a=0;a<((SMB/ZPB)/TPL);a++){ int rel_pos=threadIdx.x*(BLOCKWIDTH/TPL) + a; aRow[threadIdx.y][rel_pos] +=4*faktor*((float)value)*(1-__sinf(phi)/WAVEWIDTH); } } for(int j = 0; j < ((SMB/ZPB)/TPL); j++){ HeightMap_d[offset+threadIdx.y*WIDTH+threadIdx.x*(BW/TPL)+j] = aRow[threadIdx.y][threadIdx.x*(BLOCKWIDTH/TPL)+j]; } }

12. Probleme • CUDA-Kernel kann unter X-Linux leider nur max. 5sek laufen, bevor er terminiert wird • der Komplexität der Aufgabe ist ein Ende gesetzt ;-) • Komischerweise klappts manchmal doch • ab und an stürzt auch die GPU ab • Bildgrößen müssen vielfache von Zweierpotenzen sein • Im Idealfall: sind Zweiterpotenzen • Ursprüngliche (naive) Implementierung hatte (im Vergleich zur Finalen Version) nur mäßige Performanz • Man muss schon manchmal etwas genauer nachdenken • Dokumentation von CUDA teilweise ungenau • z.B. Shared Memory kann nicht voll ausgenutzt werden

13. Benchmarks - CPU • CPU-Implementierung ( P4 – 3,0GHz )

14. Benchmarks - CUDA • CUDA-Implementierung (GF 8800 GTX )

15. Zusammenfassung • Auf der CPU teilweise nicht zumutbare Ausführungszeiten • Allerdings noch größere Problemgrößen lösbar als mit CUDA, da kein Timeout • Primitive CUDA-Implementierung • Relativ schnell lauffähig • Speedup bereits zwischen 10 und 40 • Endversion: Enormer Speedup von ~ 500 • Bereits bei kleinen Eingabedaten • Sehr gut skalierend • Allerdings auch nur durch viel Arbeit erreichbar

16. Demo Genug geredet, jetzt wird’s gezeigt! Oder gibt’s bisher schon Fragen?