GPGPU A grafikus hardver általános célú felhasználása

1. GPGPUA grafikus hardver általános célú felhasználása Magdics Milán gumi@inf.elte.hu http://gordius.(l.)iit.bme.hu/~gumi/gpgpu

2. Tartalom • Valósidejű 3D grafika és a grafikus hardver • A mai hardverek teljesítménye • Általános célú felhasználás: • GPGPU és a grafikus pipeline • CUDA • GPGPU grafikus szerelőszalag nélkül

3. Valósidejű 3D grafika • A feladat egy (3D-s) virtuális színtér lefényképezése

4. Valósidejű 3D grafika • A cél: minél realisztikusabb kép előállítása

5. Valósidejű 3D grafika a CPU-n? • Valósidő → legalább 25 FPS • Pl.: 1280x1024-es képernyőfelbontás = 1.3M pixel →≈33M pixel / sec • Pl.: 2GHz CPU: 2*109 órajel / sec → ≈80 órajel / pixel • Ráadásul még egyéb dolgokat is számítani kellene… (fizika, AI, stb.)

6. Megjelenítési csővezeték Perspektív transzf + Vágás Virtuális világ Kamera transzformáció, 2. 1. Raszterizáció + interpoláció + textúrázás mélység szín

7. Vertex Shader Fragment Shader Grafikus kártyák Interfész Transzformáció+ Illumináció 2001-2002 Stream output Geometry Shader 2006 Vágás + Nézeti transzf + Raszterizáció+ interpoláció Textúrázás 2002 Textúra memória Kompozitálás (Z-buffer, átlátszóság) Bufferek: szín, z, stencil, …

8. Programozhatóság • A hardveren ún. shader (árnyaló) programok futtathatók • Magas szintű árnyaló nyelvek • Az utasításkészletet a grafika határozta meg • Unified Shader Model • Unified Shader Architecture Vertex Shader Geometry Shader Fragment Shader

9. Programozhatóság Visszaverődés Level of detail Füst Árnyalás Árnyékok

10. Programozhatóság HDRI Kausztikus hatások Mélység-élesség

11. A mai kártyák teljesítménye • ≈10.000.000 csúcspont/sec • ≈10.000.000.000 pixel/sec • >1 TFlop/s • Magyarország legnagyobb teljesítményű szuperszámítógépe <1 TFlop/s • SZTAKI desktopgrid csúcsteljesítmény: 3 TFlop/s • 10-100K HUF • Tesco gazdaságos szuperszámítógép • (256-1024MByte RAM)

12. A teljesítmény fejlődése

13. A teljesítmény fejlődése

14. Miért gyorsabb a GPU? • Adatfolyamok nagy számítási igényű, párhuzamos feldolgozására hozták létre • Többmagos hardver • Nagyobb tranzisztorszám jut az adatok feldolgozására – régebben csak SIMD

15. Elágazások • SIMD: minden adaton ugyanaz az utasítás hajtódik végre • Elágazásokat és változó hosszúságú ciklusokat írhatunk, de: • Ha a párhuzamos szálak más ágakon futnak, a teljesítmény jelentősen romlik • Változó hosszúságú ciklusok esetén mindig a leglassabb számít, a többi szál bevárja

16. Adat és feladatszintű párhuzamosság • Feladatszintű párhuzamosság: • Független folyamatok különböző céllal, minimális kommunikációval • Adatszintű párhuzamosság: • Sok adat, mindegyiken ugyanazt a műveletet hajtjuk végre, (általában) a többi adattól függetlenül

17. CPU vs. GPU • CPU • Gyors cache • Sok különböző szál hatékony kezelése • Egy adott szál végrehajtása hatékony Feladatszintű párhuzamosság • GPU • Sok ALU egység • Egy adott program végrehajtása minden csúcsponton/pixelen Adatszintű párhuzamosság

18. GPGPU • General-Purpose computing on Graphics Processing Units • A GPU felhasználása általános célú számításokra • Mire érdemes (mire lehet) használni? • Nagy (aritmetikai) számításigényű feladatok • Stream processing

19. Az algoritmusok leképezése GPU-ra • Adatokat kétféle primitívre képezhetjük le: • Pixelek – textúra (2D tömb) • Feldolgozás a pixel shaderben • Csúcspontok/geometriai primitívek • Feldolgozás a vertex/geometry shaderben • Megj.: GPU-CPU kommunikáció költséges

20. Adatok feldolgozása pixel shaderben • Adatok a textúrában, minden pixelre szeretnénk végrehajtani a kernelt • Rajzoljunk egy teljes képernyőt kitöltő téglalapot • Az adatok egy részének feldolgozása: több téglalap

21. Működés Kép Képpont árnyaló Számító kernel Raszterizáció Textúra Kitöltendő pixelek Minden pixel 0 Fényforrás Háromszög a képernyőn Képet lefedő négyszög Kamera 0 Csúcspont árnyaló 0 Csúcs pozíció, normál 4 db csúcs

22. Gather, Scatter • A pixel shaderben olvashatunk indirekt címzéssel (gather – x=a[i]), viszont indirekt írásra (scatter – a[i]=x) nincs lehetőség, mivel mindig a kapott pixelpozícióra írunk • A vertex shader módosít(hat)ja a pozíciót, ezáltal a kernelben határozhatjuk meg az eredmény helyét (pl. rendezés) • Megj.: az adatokat ettől függetlenül pixelként reprezentáljuk • Fullscreen quad helyett minden indexnek megfelelően egy-egy csúcspontot küldünk a GPU-ra, ezek határozzák meg a feldolgozandó adat helyét • Változó számú kimenetet is előállíthatunk (geometry shader)

23. Geometry shader • Geometry shaderrel változó számú kimenetet is előállíthatunk, ez lehet több, mint egy adat

24. Geometry shader • de lehet akár nulla is → szétválogatás, adatok kizárása a további számításokból

25. Alkalmazás: Procedurális geometria • A kirajzolandó geometriát az azt előállító algoritmussal adjuk meg • Jóval tömörebb: ~100 Kbyte vs 100 Mbyte 3D internet? • Automatikus modellgenerálás: a modell előállításához nincs szükség egy egész sereg modellező többhónapos munkájára

26. Alkalmazás: Procedurális geometria • Az algoritmikus leírás alapján a geometriát valósidőben is kibonthatjuk: • Csak azt generáljuk le, amit éppen látunk • Potenciálisan végtelen színterek (pl. egy egész univerzum)

27. Alkalmazás: Procedurális geometria

28. Korlátok • Kisméretű input, output adatok • Nincs olyan memória, ami egyszerre: • Szabadon címezhető • Írható • Globális • Nincs egyszerre írható és olvasható memória • A szálak közt nincs kommunikáció

29. GPGPU grafikus szerelőszalag nélkül • Egyszerűbben nem lehetne?

30. CUDA • Compute Unified Device Architecture • Az NVIDIA párhuzamos architektúra modellje a saját hardvereihez • GPU: sokszálú, sokmagos processzor • szálcsoportok, közös memória, szinkronizáció • Kiterjesztés a C(++) nyelvhez, a kártya programozásához • Saját fordító (nvcc), GPU kódot is generál • Könyvtárak, meghajtó és runtime

31. CUDA Alkalmazás CUDA könyvtárak CUDA runtime CUDA eszközmeghajtó Eszköz

32. CUDA • Teljesen általános célú párhuzamos számításokhoz • Nem jelenik meg a grafikus szerelőszalag (!) • Szálak, blokkok, szinkronizáció stb.

33. CUDA • A program működésének absztrakciója: • Minél több szál egy blokkban, annál jobb • De legfeljebb 512

34. CUDA • Memóriahasználat • Közös memória 16 Kbyte

35. CUDA • Működés a GPU-n: • Újabban 16 ≤ N (≤ 64) • M = 8 • Az ütemezésnek köszönhetően 4 órajelciklus alatt 32 szál fut le • 8-16K db regiszter, 64KB konstans, 6-8 KB tex. cache • Processzor ~1000-1600 MHz • Az újabb kártyákban double is

36. Korlátok a CPU-hoz képest • Kevesebb központi memória, jóval kisebb cache • Az elágazások jelentősen leronthatják a teljesítményt • Textúra: van cache, de nem írható • Az eszközmemória írható, az elérése lassú, nem cache-elt • Általában ezen lehet a legtöbbet optimalizálni • Olvasás a szálak osztott memóriájába, feldolgozás, eredmény kiírása az osztott memóriába, visszamásolás az eszközmemóriába

37. CUDA alkalmazások • Mátrix/vektorműveletek nagyméretű és/vagy sok adatra • CUBLAS könyvtár – Basic Linear Algebra Subprograms • FFT

38. CUDA alkalmazások • Sztochasztikus, Monte-Carlo szimulációk • Tőzsdei előrejelzés $ • Kockázatanalízis • Sugárkövetés

39. CUDA alkalmazások • Jelfeldolgozás, szűrések, képfeldolgozás

40. CUDA alkalmazások • Fizikai szimulációk

41. CUDA alkalmazások • PET rekonstrukció

42. CUDA alkalmazások • PET rekonstrukció • Problémák: • 35x35 detektor / modul (lap), mindegyik 3x35x35 másikkal van koincidencia kapcsolatban ~ 3 millió pár • Cél: 512x512x512=227 felbontású adat • Minden párba érkezhet foton a térfogat minden pontjából

43. Programozzunk! • Szerezzünk NVIDIA kártyát  • http://www.nvidia.com/object/cuda_home.html • CUDA Driver, CUDA Toolkit, CUDA SDK (sok példakóddal)

44. Programozzunk! • CUDA kernel: ami a GPU-n egy szál lesz • ~ shader program • C függvény // kernel definíció __global__ voidvecAdd(float *A, float *B, float *C) { C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; }

45. Programozzunk! • Kernel hívás: int main() { // kernel hívás vecAdd<<<1,N>>>(A ,B,C); // … } Grid dimenzió Blokk dimenzió

46. Programozzunk! • Erőforrások: • Pointer jelenthet egy központi (CPU) memóriabeli, és egy eszköz (GPU) memóriabeli címet is, a kettő nem összetévesztendő // CPU memória h_p = malloc(sizeof(int)*42); // GPU memória cudaMalloc((void**)&d_p, sizeof(int)*42); // másolás a kártyára cudaMemcpy(d_p, h_p,sizeof(int)*42, cudaMemcpyHostToDevice);

47. Programozzunk! • Szinkronizációs utasítások: • Szálakon belül: az egy blokkban futó szálak bevárják egymást • Host kódban: a CPU futása felfüggesztődik, amíg le nem fut minden kernel (a kernel indítás aszinkron)

48. Jó tudni • Debugger még nincs, de van emulációs mód • Minden a CPU-n fut • Host kód a kernelekben • Shader műveletek és típusok: <cutil_math.h> • A kernelek futási idejét az op. rendszer limitálja • Registry „hack”, CUDA ablakozó rendszer nélkül • Vagy: szedjük szét rövidebb kernelekre #ifdef __DEVICE_EMULATION_ // hoston futó kód, csak emuláció esetén fut le

49. Házi feladat  • 1. Kínában matematika érettségi zajlik. A feladatsor tesztes, minden feladatra 4 válaszlehetőség van, a teszt 32 kérdésből áll. Az érettségin 10(-100) millió diák vett részt. Rendelkezésünkre állnak a kérdésekre adott válaszaik, valamint a megoldókulcs. Hogyan állapítsuk meg az átmenő jegy (2?) ponthatárát, hogy a diákok kb. 10%-a bukjon meg?

50. Házi feladat  • 2.a. Írjunk isosurface megjelenítőt. Adottak egy térfogat (3D tömb) sűrűségértékei. Jelenítsük meg a térfogat egy adott (állítható) sűrűségű részeit. Tipp: http://en.wikipedia.org/wiki/Marching_cubes • Vagy: • 2.b. Jelenítsük meg a térfogatnak a pixelben látható maximális sűrűségét.