GPU-Computing mit CUDA und OpenCL in der Praxis

  • Published on
    07-Jul-2015

  • View
    1.584

  • Download
    4

Embed Size (px)

DESCRIPTION

Folien des Vortrags auf der Parallel 2012 am 24.5.2012 in Karlsruhe.

Transcript

  • 1. GPU-Computingmit CUDA und OpenCL in der Praxis 24. Mai 2012Jrn Dinkla

2. Motivation 3. Computer-Grafik Vom Modell zum BildTomas Akenine-Mller 2002, Quelle: http://www.realtimerendering.com/ 4. Echtzeit-GrafikVertexPixel Quelle: Tomas Akenine-Mller 2002 5. Rckblick1992 1995 6. T&L in HardwareFixed pipelinesDirectX 7.019992000 7. Programmierbare ShaderDirectX 8.0 DirectX 9.020012002 8. Shader-SprachenGPGPUHLSL, Cg, GLSL2003 2004 2005 9. Unified ShaderDirectX 10.020062007 10. GPU-Computing CUDAOpenCL CTM Stream SDKDirectX 11.0200620072008 2009 11. Massiv Parallel 12. Parallelisieren? Schneller, Grer, Besser480p 576p720p 1080p 4K2D / QFHD Hhere Geschwindigkeit 13. Entwicklungs-Methodik1. Golden Code Sequentiell und korrekt2. Golden Code parallelisieren3. Zu GPU-Code transformieren 14. Parallelitt findenEingabeEingabeVerarbeitung#1 #2#3 Ausgabe Ausgabe 15. Parallelitt findenEingabeE1 E2 E3Verarbeitung #1 #2 #3 Ausgabe A1 A2 A3 16. SPMD / SIMT Parameter: Eindeutige ID Hole Daten anhand ID Verarbeitung Speichere Daten anhand ID E1 E2 E3 #1 #2 #3 A1 A2 A3 17. SIMT: Lock-step1 23 41 23 4 uchar4 p; p.x = x+y; p.y = x-y; p.z = y; p.w = 255; 18. FrameworksPlattform-Unabhngigkeit ?Komfortabel OpenACC ? C++ AMP?PyCUDA /BesserThrust? PyOpenCL ?CUDA Runtime PraktischAPI C++ cl.hppLow Level CUDA Driver APIOpen CL Hardware 19. CUDA Runtime API HelloWorld.cuKernelHost 20. Programm zu Code Host CodeDevice CodeC/C++ (Kernel) JIT?CompilerCompiler Assembler / LLVM / PTX / IL LLVM JIT? AssemblerAssembler (*) Cubin / MachineMachine CodeCode 21. OpenCL Vorteile Plattformunabhngig CPU+GPU Vektorberechnungen Aber Performance von Device abhngig Komiteesprache 22. OpenCL / Driver API PlattformDevice ContextCommand Queue ProgramKernelAufruf 23. WebCLSiehe http://webcl.nokiaresearch.com/kerneltoy/ 24. Smoothing 25. Smoothing 3x3 Fenster0 1 2 3 4 0 1 2 3 4 0 1 2 3 40 0 01 1 12 2 23 3 34 4 4 26. Algorithmus Fr alle x,y,z in vol sum = 0, c = 0 Fr alle dx,dy,dz in Nachbarschaft sum += vol[x+dx,y+dy,z+dz] c += 1 vol[x,y,z] = sum / c Threads 27. Extent Extension / Gre width, height, depth index(x,y,z) inBounds(x,y,z)0 1 2 300 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15123 28. Accum Akkumulator0 1 2 3 4 add(value) 0 Akkumuliert int4 1 avg() 2 3 Durchschnitt 4 29. Golden Code 30. Fermi (GTX 580) 31. Kernel-Konfiguration x*y*z Threads Cores/SMs Thread-Block / Work group / Tile Grid / NDRange0 1 2 3 4 5 6 70 T T T T T T T T1 T T T T T T T T2 T T T T T T T T3 T T T T T T T T4 T T T T T T T T5 T T T T T T T T6 T T T T T T T T7 T T T T T T T T 32. Kernel-Konfiguration ExecConfig grid threads stream ExecConfig(Extent) 33. Hostberblick Buffer AlgorithmusBufferBufferDeviceBufferKernelBuffer 34. Speicher-Management Device-Speicher cudaMalloc(), cudaFree() Host-Speicher cudaMallocHost(), cudaFreeHost() Transfer cudaMemcpy() 35. Host-SpeicherVirtuellerSpeicher Physikalischer Speicher cudaMallocHost() malloc() DeviceSpeicher 36. Buffer BaseBuffer malloc(), free() HostBuffer PinnedHostBuffer DeviceBuffer copyFrom(), copyTo() Versionierung 37. BufferPair Paar HostBuffer Buffer Host-Buffer Device Buffer Kernel Buffer Device-Buffer Methoden updateDevice() updateHost() 38. GPU-Code 39. Performance-Vergleich Nur GPUMit KopienGreCPUGPU Speedup GPU+C Speedup800 20,0016 10 20,0032 20 40,0064150 40,00 12897424,2510 9,70 2566602328,706310,48 384 22167828,41 20410,86 512 5249 184 28,5348210,89Und grer ? 40. Speicherplatz: N^3 41. Grere Volumen SwappingHostDevice 42. Swapping BufferIterator Kernel: Anpassen BufferPair: Erweitern 43. Host Auslastung I O I O I ODevice I K O I K O I K OLast 44. Streams & OverlapSchritt Ein Stream Stream 1Stream 2 Stream 1Stream 21 H2D H2DH2D2 Kernel 1Kernel 1H2DKernel 1H2D3 D2H D2HKernel 2D2HKernel 24 H2D D2HH2D D2H5 Kernel 2H2DKernel 36 D2HKernel 3D2H7 H2D D2H8 Kernel 39 D2H Kernel + Kopie Kernel + Kopie undberlappend H2D und D2H 45. Swapping & Streaming Initialisierung Parallel Fr alle ungeraden Partitionen p Kopiere H2D fr p in s1 Rufe Kernel auf fr p in s1Asynchron Kopiere D2H fr p in s1 Fr alle geraden Partitionen q Kopiere H2D fr q in s2 Rufe Kernel auf fr q in s2 Kopiere D2H fr q in s2 46. bersicht Asynchrone Kopien HostBufferBuffer Buf 1 Kernel Buf 1 Device Buf 2Kernel Buf 2 47. Double Buffering Volumen VolumenInput OutputHost Buf 1 Buf 2 Out 1 Out 2 Buf 1 Kernel Out 1Device Buf 2Kernel Out 2 48. Klasse HostBufferPair Analog zu BufferPair Buffer HostBuffer PinnedHostBuffer Unterschiedliche Gre updateFrom(), updateTo() 49. Performance-Vergleich GreCPUGPUSpeedup 8 00 1600 3210 0 6420 0128 110 0256 660 23 10,48384 223380 10,86512 5263 183 10,8976817707 171810,311024 42101 407910,321152 59156 59249,99 50. Theorie GTX 580 1632,3 GFLOPS und 194,5 GB/s Wen es interessiert: GFLOPS #Cores * Shader-Takt * 2 GB/s Breite [Byte] * Memory-Takt * x GDDR3: x = 2, GDDR5: x = 4 51. Pi mal Daumen Pro Kernel 27 Laden, 1 Speichern 638 Operationen Arithmetische Intensitt Bei 512er Test 467,93 GFLOPS 20,54 GB/s Optimierungspotential vorhanden! 52. OptimierungMaximiere 1. Parallelitt2. Speicherdurchsatz3. Berechnungsdurchsatz 53. Maximiere Parallelitt Streams Asynchrone Kopien und Kernel Auslastung des Device Grid Auslastung der SMs Thread-Block Occupancy 54. Occupancy Mehrere Thread-Blcke pro SM Speicherzugriffe verstecken O(1) Scheduler SM TB TB TB TB TB -- -- -- 55. Occupancy Calculator 56. Speicher-Architektur GPUGlobal Memory Global Memory8-20 8-16Constant Texture Bus / Memory Controller 160-L2 Cache 200 CPU 1Prozessor (SM) CoreCoreC C CC 8000 L1 / L2 L1 / L2 RegistersL3 Cache Local / Shared / L11600 57. Max. Speicherdurchsatz Coalesced Zugriffe 32, 64, oder 128-byte Alignment 58. Max. Speicherdurchsatz Pinned-Host-Speicher Minimiere Kopien Benutze On-Chip-Speicher Shared/Local, Register Konfiguriere Cache 59. Divergenz1 2 341 2 34 int tid = treadIdx.x; if (tid < 2) {o[tid] = 1; } else {o[tid] = 2; } 60. Max. Berechnungen Minimiere Divergenz Loop-Unrolling Berechnen statt Speichern Arithmetik Przision vs. Geschwindigkeit Fusion von Kerneln 61. Synchronisation Innerhalb Thread-Block __syncthreads() Atomare Zugriffe atomicAdd() Speicher Zwischen Kernel-Aufrufen 62. Weiterfhrend 63. JVM JNI JCuda Javacl, Jocl Eigenes API Aparapi, Java-GPU 64. JVMKomfortabelBesserCUDA Runtime PraktischAPI C++ cl.hpp JCUDALow Level CUDA Driver API Open CL JavaCL Hardware 65. CUDA mit Java 66. Hello, Groovy CUDA! JCUDA 67. Hello, Groovy OpenCL!JavaCL 68. Hello, Scala + GPU JavaCL 69. Fazit JVM Vorteile Entwicklungsgeschwindigkeit Host-Code Nachteile Datentypen Getrennt Debuggen 70. Fazit Richtig eingesetzt unschlagbar! Folien & Code http://dinkla.net/parallel2012 71. Literatur: CUDA Sanders, KandrotCUDA by Example Kirk, HwuProgramming Massively Parallel Processors 72. Literatur: OpenCL ScarpinoOpenCL in Action Gaster et.al.Heterogeneous ComputingWith OpenCL 73. Literatur: CUDA Hwu (Ed.)GPU Computing GemsEmerald Edition Hwu (Ed.)GPU Computing GemsJade Edition

Recommended

View more >