540 likes | 661 Views
GPGPU A grafikus hardver általános célú felhasználása. Magdics Milán gumi@inf.elte.hu http://gordius.(l.)iit.bme.hu/~gumi/gpgpu. 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
E N D
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
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
Valósidejű 3D grafika • A feladat egy (3D-s) virtuális színtér lefényképezése
Valósidejű 3D grafika • A cél: minél realisztikusabb kép előállítása
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.)
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
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, …
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
Programozhatóság Visszaverődés Level of detail Füst Árnyalás Árnyékok
Programozhatóság HDRI Kausztikus hatások Mélység-élesség
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)
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
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
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
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
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
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
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
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
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)
Geometry shader • Geometry shaderrel változó számú kimenetet is előállíthatunk, ez lehet több, mint egy adat
Geometry shader • de lehet akár nulla is → szétválogatás, adatok kizárása a további számításokból
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
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)
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ó
GPGPU grafikus szerelőszalag nélkül • Egyszerűbben nem lehetne?
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
CUDA Alkalmazás CUDA könyvtárak CUDA runtime CUDA eszközmeghajtó Eszköz
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.
CUDA • A program működésének absztrakciója: • Minél több szál egy blokkban, annál jobb • De legfeljebb 512
CUDA • Memóriahasználat • Közös memória 16 Kbyte
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
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
CUDA alkalmazások • Mátrix/vektorműveletek nagyméretű és/vagy sok adatra • CUBLAS könyvtár – Basic Linear Algebra Subprograms • FFT
CUDA alkalmazások • Sztochasztikus, Monte-Carlo szimulációk • Tőzsdei előrejelzés $ • Kockázatanalízis • Sugárkövetés
CUDA alkalmazások • Jelfeldolgozás, szűrések, képfeldolgozás
CUDA alkalmazások • Fizikai szimulációk
CUDA alkalmazások • PET rekonstrukció
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
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)
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]; }
Programozzunk! • Kernel hívás: int main() { // kernel hívás vecAdd<<<1,N>>>(A ,B,C); // … } Grid dimenzió Blokk dimenzió
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);
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)
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
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?
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.