Párhuzamosítás Amikor egy szál nem elég...
Párhuzamosítás • Algoritmikus optimalizáció > párhuzamosítás – Soha sincs végtelen feldolgozóegység
• Legjobb esetben lineáris gyorsulás – Függőségek, közös erőforrások (memória sávszélesség, hálózat, ...) miatt általában ennél kevesebb
• Mekkora lineáris gyorsulás? – 1 szál helyett akár 4096 16 szálat futtató mag = 65536szoros gyorsulás • 1 nap = 86400s azaz ami egy napig futott az akár 1s lehet • Annyira mégse rossz 2
Párhuzamosítás • Van órajel is, ami viszont általában fele vagy negyede – Így se rossz az 5 másodperc egy nap helyett
• Mikor lehet párhuzamosítani? – Nem triviális – Akkor biztosan nem lehet, ha mindig meg kell várni az előző lépés eredményét. – Képeknél általában jól lehet, mert egy kimeneti pixel csak a bemeneti képtől szokott függeni 3
Műveletek képeknél • Pixel manipulációk (két kép összeadása, intenzitás szorzása,...) – Triviálisan nagyon jól párhuzamosítható
• Minimum, maximum, összeg, normalizálás az egész képre, FFT – Parallel reduction
• Szűrések – Minden pixel a saját környezettől függ – A triviális megoldás is gyorsít • Szeparábilis szűrők 4
Műveletek képeknél • Szűrések (folytatás) – GPU-n érdemes a nem triviális megoldást választani – Nem konvolúciós szűrők • Mediánnál pl. a rendezés nem megy egy lépésben • A morfológia is ilyen! De az könnyen számolható
• Szegmentáció, flood fill – A már meglévő eredményt felhasználjuk a további iterációkban. • Azért itt is jelentősen lehet gyorsítani 5
Műveletek képeknél • Rekonstrukció, Röntgen szimuláció, Vizualizáció – „Mindenki mindenkitől” függ kis túlzással – Több iteráció – Egy iteráción belül a kimeneti pixelek külön számolhatóak
• Optical flow, regisztráció – Ezek se lokális algoritmusok
• Hisztogram – Sok száll írja a kimenetet -> atomi műveletek 6
Párhuzamosság • Statikus párhuzamosság – Az elején eldöntjük, hogy hány szál fusson – Pl. két kép összeadása – Aktív szálak száma így is változhat • Flood fill
• Dinamikus párhuzamosság – Menet közben növekedhet a szálak száma – Pl. ray tracing 7
Megkötések • Nem minden eszköz tud dinamikus párhuzamosságot • Lehetnek feldolgozóegységek amelyek nem képesek független műveletet végrehajtani – SIMD • MMX, SSE, AVX • GPU Multiprocesszorok
• Közös erőforrások elosztása – Leggyakrabban a memória sávszélesség a kritikus 8
Megkötések • SMP vs NUMA – SMP - Symmetric multiprocessing • Minden feldolgozóegység ugyanúgy fér hozzá az egész memóriához • GPU ilyen – De itt is jelentősen befolyásolja a teljesítményt, hogy az egymást követő szálak egymást követő memóriacellákat érnek-e el.
– NUMA - Non-uniform Memory Access • Új Intel CPU-k ilyenek, a magokhoz tartoznak memóriák • Többprocesszoros rendszereknél akár igen nagy költsége is lehet egy távolabbi egységhez tartozó memória elérésének 9
Hátrányok • Olvashatatlan lesz a kód ha nagyon optimalizáljuk • Nehéz debuggolni – Szerencsére azért már vannak eszközök GPU-ra is de nem olyan jók, mint a CPU-sak. – Több szálat CPU-n is nehezebb debuggolni
• Új nyelveket, API-kat kell megtanulni • Nem hordozható a kód – Vagy legalábbis elveszítjük az előnyöket ha nagyon optimalizálva volt 10
Eszközök • CPU – Több szál indítása – OpenMP – Vektorizáció • Intel: MMX, SSE (2,3,SSSE3,4.1,4.2), AVX, AVX2 – Xeon: AVX-512
• ARM: Neon
• GPU – OpenCL, CUDA, GLSL/HLSL
• Gyorsító kártyák – Intel Xeon Phi – FPGA
• Szabad több gépet is használni 11
CPU • Előnyök – – – –
Az egész memóriát elérjük és sok memóriánk van Akár 64Gb memória hagyományos asztali gépben Jól debuggolható Könnyen kezelhető a már valószínűleg ismert eszköztárral – Könnyű szinkronizáció – Magas teljesítmény szálanként – Az adatokat nem kell fel-/letölteni
• Hátrányok – Nem annyira párhuzamos, mint a többi megoldás 12
GPU • Előnyök – Nagyon párhuzamos – Elérhető ár, jó ár/teljesítmény arány – Jó FLOPS/Watt arány
• Hátrányok – Külön nyelvet kell megtanulni – Nagyon specifikus optimalizáció – Külön memória, fel-/letöltés 13
Gyorsító kártyák • Előnyök – Speciális feladatokra lehet optimalizálni a hardwaret
• Hátrányok – Drága – Specializált, tehát nem hordozható és meg kell tanulni hozzá mindent
14
CPU Technológiák • CPU – Processek közötti kommunikáció – _beginthread() / fork() • Mutexek • Események
– OpenMP – Auto-vektorizáció – Intrinsic
15
Processzek közötti kommunikáció • Futtathatunk párhuzamosan több programot – Windowson és Linuxon is van lehetőség közös memóriaterület létrehozására – Lehet rendszer szintű mutexekkel, eseményekkel szinkronizálni
• Hasznos lehet két külön fejlesztett modul összekötésénél, több különböző kimenet előállításánál • Pl. CAD szerver, 3D rekonstrukció, megjelenítő párhuzamosan futhat • Az erőforrásokra nehezebb így figyelni 16
Szálak indítása • Windows – _beginthread()/_beginthreadex() • Windows.h
– Egy függvény a paramétere aminek átad egy void * paramétert – Viszonylag nagy overhead az elindítás – Az ütemezés szálanként megy, így több processzoridőt kapunk eleve – Szinkronizáció • Mutex – Csak egy szál birtokolhatja – Akár rendszer szintű
• Critical section
17
Szálak indítása • Windows – Szinkronizáció (folytatás) • Események – Lehet várakozni rájuk – Lehet triggerelni őket
• Linux/Unix – fork() • A két szálon más a visszatérési értéke
– Nagyjából ugyan azok a szinkronizációs eszközök állnak rendelkezésre, mint Windowson 18
Szálak indítása • Jó ugyanarra, mint a több processz indítása, csak nehezebb cserélni egy részét • Hasznos ha egy folyamatnak több részét tudjuk párhuzamosan csinálni – Pl. Pipeline program szervezés • Következő kép betöltése/előfeldolgozása miközben az aktuálisat feldolgozzuk
19
OpenMP • C/C++ #pragma • Thread pool-okat használ – Sokkal kisebb overhead egy szál indítása
• Nagyon kis módosítás a programkódban • Párhuzamos for ciklusok #pragma omp parallel for for(int i = 0;i < 10000;i++) { A[i] = i; } 20
OpenMP • Be kell kapcsolni a fordítónál – gcc: -fopenmp – Visual Studio: C/C++ / Language / OpenMP support
• omp.h – Függvények amivel a pragmák mellett megadhatunk paramétereket • pl. A szálak számát
21
OpenMP blokkok • Parallel blokkok – Egy paragma egy párhuzamos blokkot hoz létre – Egy blokkon belül nem lehet még egy párhuzamos blokk • OpenMP 3.0-tól meg lehet oldani
– Több egymásba ágyazott for párhuzamosítása egy for ciklusra való átírással
22
OpenMP változók • A szekción belül definiált változók privátok – Minden szálnak saját példány
• A szekción kívül definiált változók shared-ek – Egy közös változó az összes threadnek
• Explicit módon is meg lehet adni: #pragma omp parallel for private(a, b) shared(c)
23
OpenMP redukciók • Redukció amikor sok adatból lesz egy – Szumma, átlag, szorzat, és kapcsolat,...
• Atomi műveletek kellenének – Elfedi előlünk az OpenMP #pragma omp parallel for reduction(+:sum)
24
OpenMP omp.h • Vannak hasznos függvények: – omp_get_thread_num() – omp_get_num_threads() – omp_set_num_threads() • Nem garantált, hogy tényleg annyi lesz
• A legtöbb esetben nem szükségesek a függvények #pragma omp parallel for num_threads(4)
• Az OpenMP még sok mindenre képes, de az esetek 90%-ban ennyi elég 25
Auto-vektorizáció • A fordítót meg lehet kérni, hogy for ciklusokat írjon át SIMD utasításokra – Gcc –O3, Visual C++ -Ox
• A fordító buta és mindig a legrosszabbra készül – __restrict__ pointerek • Biztosítja, hogy csak az a pointer és annak másolata éri el a megcímzett területet • Még nem C++ szabvány, de C99 és szinte mindenki támogatja 26
Auto-vektorizáció • A fordító nem tud semmit az alignmentről se – gcc: __attribute__(aligned(16)) – Visual C++: __declspec(align(16))
• Csak 4,8,16 elemű for ciklusokat vektorizál – Ami pont egy utasítás
• Nem használja ki a maszkolás lehetőségét se • A gcc buta, a Visual C++ még butább – Megoldás: Intrinsic-ek 27
Intrinsic • Függvényeknek álcázott assembly utasítások – Nem a legolvashatóbb kódot kapjuk vele
• A fordító foglalkozik a regiszterekkel – Mi használhatunk változóneveket
• Egy intrinsic függvény garantáltan egy utasítássá fordul • Be kell kapcsolni a megfelelő architektúra utasításait – gcc –m<architektúra> (pl.: -msse2) – Visual C++: C/C++ / Code Generation / Enable Enhanced Instruction Set / <architektúra> 28
Intrinsic MMX • MMX – MM0-MM7 64bites regiszterek • Az FPU regiszterinek az alsó 64bitje
– Egyszerű SIMD utasítások egész számokhoz • 2x 32 bit, 4x 16bit vagy 8x 8bit
– Pentium with MMX és Pentium II óta minden támogatja • (ezek 2000 előtti processzorok)
– Nem igazán használjuk már, van jobb 29
Intrinsic SSE • XMM0-XMM7 128bites regiszterek – 64biten további 8 regiszter: XMM8-XMM15
• Floating point SIMD utasítások – 4x 32bit float műveletek
• Gyorsabb a memória elérés ha 16 bájtos határra vannak igazítva a címek • Pentium III óta minden támogatja (1999) • (Oprendszer támogatás kell hozzá) 30
Intrinsic SSE2 • Innentől kezd igazán hatékony eszköz lenni • Az XMM regiszterekben több típus használható – 2x 64 bit double – 4x 32 bit float – 4x 32 bit int, 8x 16bit short, 16x 8bit char
• MMX utasítások az XMM regiszterekre • Pontatlanabb, mint az FPU (64bit vs 80bit) 31
Intrinsic SSE3 • Horizontális műveletek az XMM regisztereken – Vertikális amikor összeadok két regiszterben minden mezőt párhuzamosan – Horizontális amikor egy regiszterben összegzem az elemeket
• Pentium 4 NetBurst mikroarchitektúra óta (2000)
32
Intrinsic SSSE3, SSE4.1, SSE4.2 • További horizontális utasítások az MM és az XMM regiszterekhez • Néhány további keverés utasítás • Bájtonkénti keverés • Minimum, maximum utasítások • Egyéb mindenféle spéci utasítások kihegyezve egy egy konkrét feladatra • Core mikroarchitektúra óta SSSE3 (2006), 2008-as Core i7 óta SSE4 33
Intrinsic AVX • Új 256 bites YMM regiszterek – Alsó 128 bitben az XMM regiszterek
• SSE utasítások kiterjesztve 256 bitesre • 4x 64bit double, 8x 32bit float • Nem destruktív SSE utasítások – Eddig: a = a+b – Mostantól: c = a+b • A 128 bites XMM regiszterekre is
• Ez már 8x gyorsítás float számításoknál! • Sandy Bridge óta (2011) elérhető 34
Intrinsic AVX2, AVX-512 • Haswell óta AVX2, még nem túl sok gép támogatja... • Integer utasítások az YMM regiszterekre • Néhány további floating point utasítás – FMA – FFT-re, skalár szorzatra optimalizálva
• AVX-512 a legújabb Xeonokban nemsokára – ZMM regiszterek 512bitesek, alsó 256 az YMM – Kiterjesztve az AVX, AVX2 utasítások – További 2x-es gyorsítás 35
Intrinsic • Intrin.h – Itt vannak a C++ deklarációk – __m256 típus az YMM regiszterekhez __m256 a,b,c; a = _mm256_set_ps(1,2,3,4,5,6,7,8); b = _mm256_set_ps(8,7,6,5,4,3,2,1); c = _mm256_add_ps(a,b); //c = [9 9 9 9 9 9 9 9];
– ARM architektúrán az arm_neon.h-ban vannak ARM-os SIMD utasítások 36
Intrinsic + OpenMP • Az OpenMP több magot fog be munkára • A vektorizáció egy processzormagon párhuzamosítja a feldolgozást • A gyorsulás mértéke a kettő együttes használatával a kettő szorzata • Core i7 esetén (2x4 mag HT-vel) az egyszálas programhoz képest floatokkal számolva akár 64-szeres gyorsulás! (Ideális esetben) 37
GPU Technológiák • GLSL/HLSL • GPGPU – OpenCL – CUDA
38
GLSL/HLSL • • • •
OpenGL 2.0-tól GLSL shader nyelv DirectX HLSL Hasonlít a két nyelv Létrehozhatunk framebuffereket (nem csak a képernyőn) és tudunk bele renderelni • Vertex Shader nem csinál semmit csak átnyomja magán a pontokat • Fragment Shader pixelenként megy, ebben tudunk érdemi feladatokat megcsinálni • Sok beépített függvény, egyszerűbb inicializáció, elég jól hordozható 39
GPGPU • OpenCL és CUDA nagyjából ugyan azt tudja • OpenCL nem csak GPU-kat támogat és nem csak Nvidia eszközökön használható • Az Nvidia ott tesz keresztbe az OpenCL-nek ahol tud – Végre támogatják az OpenCL 1.2-őt (2.1-nél járunk)
• A CUDÁnak jobb a szoftvertámogatottsága 40
GPGPU: CUDA, OpenCL • Közös vonások • Adatok fel-/letöltése – PCIExpress busz overheadje
• Kernel hívás – Viszonylag nagy overhead
• Kernelhíváson belül nincs globális szinkronizáció! – Új kernel hívásával pedig nagyon költséges 41
GPGPU: OpenCL, CUDA • CUDA terminológia szerint egy kernelhívásnál van egy Grid • A Grid tartalmaz blokkokat • A blokkok futtatnak szálakat • Van globális és lokális id minden szálhoz – Globális a Grid-en belüli hely – Lokális a blokkon belüli hely
42
GPGPU Memória • Globális memória – A kártyára integrált GDDR5 memória – Sok (2-8Gb)
• Lokális memória (Shared memory) – A GPU belsejében a multiprocesszorhoz tartozó memória – Egy blokk ezt látja
• Privát memória – Egy szál látja (regiszterek kb.) 43
GPGPU Memória • Konstans memória – Mint a globális memória, de úgy van cachelve, hogy minden szál könnyen tudja párhuzamosan olvasni
• Textúra memória – Mint a globális memória, de a textrázó egységek tudnak benne interpolálni és ehhez is van külön cache
44
GPGPU Memória • Privát memória – 1 órajel ciklus
• Lokális memória – 10 órajel ciklus
• Globális memória – 100 órajel ciklus
• Konstans és textúra memória – 100 órajel ciklus amíg nincs cachelve, utána 10 45
GPGPU Utasítások • Gyakorlatilag nincs ugró utasítás – Az if mindkét ága lefut, de maszkolni lehet, hogy melyik threadekben hajtódjon ténylegesen végre – Nem jó a nagy elágazási tényező • Fmax, fmin használata sokszor segít • Nem hatékony ha egy blokkon belül a szála különböző ágára futnak az if-nek
• Nem túl hatékonyak a nagy for ciklusok – Fejtsük ki az összes ciklusát egy for-nak (ha rövid) • #pragma unroll
• A double 16-szor lassabban számolódik, mint a float 46
GPGPU Memória elérés • Amit többször is használunk töltsünk be a lokális memóriába – Ha csak az az egy szál használja, akkor a privát memóriába
• A ténylegesen párhuzamosan futó szálak számát befolyásolja a privát memória használata – Kb. 1kb ami egy szálnak jut ideális esetben 47
GPGPU Memória elérés • A memóriából csak sorban lehet olvasni és csak 256 bájtos határra igazított blokkokat – Ha nincs igazítva az adat és átcsúszunk akkor két olvasás történik – Ha nem sorban olvasunk akkor több olvasás történik • Egyre több mintát támogat a memória vezérő, nem csak a tökéletesen rendezett sort
• A memória egy bankját egyszerre egy blokk érheti el, ha több blokk hivatkozik rá, akkor meg kell várniuk egymást – Konstans memória használata olvasáskor segít 48
GPGPU Kernelhívás • A Grid méret mindig osztható kell, hogy legyen a blokk mérettel • Érdemes egyetlen kernelhívást csinálni és mindent abban kiszámolni • Az adatok fel-/letöltése miatt csak nagy problémák esetén érdemes GPU-t használni • Minél több szálat indítunk annál jobb a GPU kihasználtsága – 1’000’000 szál még nem olyan sok, hogy tökéletesen kihasználja a GPU-t 49
GPGPU Cache • A textúra úgy van cachelve, hogy 2D-ben a közeli pontok is betöltődjenek egy olvasáskor • CUDA esetén a shared memory és a cache közös tárolóban van – Ugyan olyan gyorsak – Érdemes előre betölteni amit használni fogunk, amennyiben többször is olvassuk
50
GPGPU textúra memória • Nem mindig írható • Textúrák speciálisan cachelve – 2D-ben Szomszédos pixelek töltődjenek be, ne sorfolytonosan • Csigavonl minta a tároláshoz
• Lineáris interpoláció hardware támogatása – Ugyan olyan gyors, mint egy pixel kiolvasása – Fel lehet használni biköbös interpolációhoz is • 4 olvasás a 16 helyett 51
Lokális szinkronizáció • Egy blokkon belül bevárhatják egymást a szálak – Hasznos például a lokális memória feltöltése után – Lokális memóriában elvégezhetnek redukciós feladatokat – Bármilyen más blokkon belüli kommunikáció
52
OpenCL • Macerásabb inicializálás – Cserébe nem csak GPU-kat támogat
• • • • •
Platform kiválasztása Eszköz kiválasztása a platformon Kontextus létrehozása Commad queue létrehozása Egy GPU esetén így kész vagyunk, több GPU esetén többet kell ezekből létrehozni 53
OpenCL kernelhívás • Kernel külön .cl fájlban forrásként mellékelve – Buildelni kell a program indulásakor -> program objektum • Lassú
– Minden kernelből kernel objektumot kell csinálni
• Minden kernel híváshoz be kell állítani minden paramétert egy-egy függvényhívással • Még egy függvény hívás elindítja (akár aszinkron módon) a kernelt – AMD aszinkron, NVidia szinkron – Itt kell megadni a globális méretet 1,2 vagy 3 dimenzióban – A lokális méret opcionális, rábízhatjuk a driverre
• Nem túl beszédesek a hibaüzenetek 54
OpenCL textúrák • 1D, 2D, 3D, 2D textúra tömbök – Textúratömbök OpenCL 1.2 óta
• 1D, 2D, 2D textúra tömbök írhatóak • 3D csak kiegészítéssel írható GPU oldalról – Tömb írható és a tömböt bele lehet másolni
• Van Nearest Neighbour, Lineáris interpoláció – Hardver támogatás, a lineáris interpoláció ugyan olyan gyorsan olvas! – Textúrázó egységek számától függ a teljesítmény 55
OpenCL kernelek __kernel void immul(float q, __write_only image2d_t dest, __read_only image2d_t src ) { sampler_t pixelsampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP_TO_EDGE| CLK_FILTER_NEAREST; int2 id = (int2)(get_global_id(0),get_global_id(1)); float tmp = read_imagef(src,pixelsampler,id).x*q; write_imagef(dest,id,(float4)(tmp,0,0,0)); } 56
OpenCL típusok • C99 szabvány szerinti skalárok • Vektor típusok – Int2, int4 – float2, float4
• Van mindennek CPU-s megfelelője – cl_float2 mezőit az s tömbön keresztül érjük el cl_float2 a; a.s[0] = 0; a.s[1] = 42;
• A vektor változók használata csak az olvashatóságot javítja, nem gyorsítanak • A vektorváltozóknál az operátorok elemenként hajtódnak végre 57
OpenCL függvények • Szintén C99 • Matematikai – fabs, fmin, fmax, sin, cos, tan,... – Vektor műveletek • dot, cross,...
• Id és méret lekérdezés • Textúra kezelés • OpenCLUtils.h és .cpp-t írtam már, printf szerű kernel hívások, inicializáció – Csak egy GPU kezelése – Felrakjuk a tárgyhonlapra
58
CUDA • Egyszerűbb inicializáció – Csak NVidia GPU-val megy
• Ha egy GPU van a gépben már tölthetjük is fel az adatokat és hívhatjuk a kernelt • .cu fájlok – Az nvcc szedi szét C++-ra és CUDA C-re – A C++ tovább megy a fordítónak – A CUDA C-ből ptx kódot csinál az nvcc 59
CUDA • Elég a ptx fájlokat adni a projekthez – Lehet forrást is adni
• A GPU-ra írt kód keveredik a CPU-s kóddal, spéci függvények a kernelek.
60
CUDA Kernelhívás • Olyan, mint egy template függvényhívás kernel_ipm_extend <<< gridsize, blocksize >>> ( gDst_ipm, dst_y, dst_x, dst_y_memstride );
61
CUDA textúrák, kernelek, típusok • • • • • • •
Mint OpenCL-nél De itt mindent lehet írni Template-re hasonlító szintaxis itt is Interpoláció itt is hasonlóan Itt a kernel egy __global__ függvény A méret és az id kitüntetett változókban Hasonló típusok itt is, ugyan úgy csak az olvashatóságot javítja a vektorizáció 62
CUDA függvények • Hasonló beépített függvények • Sok jó könyvtár van hozzá – CuFFT – CuBLAS
• Az OpenCL-hez is van FFT, BLAS, meg pár könyvtár – A dokumentáció ott rosszabb, és körülményesebb használni 63
GPGPU Esettanulmányok • Szűrés – Naiv – Lokális memória használata
• Flood fill – Naiv – Lokális memória használata – Ping-pongozás
• Parallel reduction – Szummázás
• Hisztogram számítás 64
GPGPU Naiv szűrés • Minden pixel párhuzamosan számolódik • Minden pixel beolvassa a számára szükséges bemeneteket • For ciklusban konvolúció – Esetleg egyéb művelet nem konvolúciós szűrőnél
• Eredmény kiírása • Jelentősen többször olvaszuk be a globális memóriából a bemeneti kép pixeleit, mint amennyiszer szükséges 65
GPGPU Szűrés • Rakjuk a lokális memóriába a pixeleket • A for ciklus futhat a lokális memórián • A blokk szélén több pixelt kell beolvasni, hogy ne lógjunk ki – Az if használata nem jó, és ilyenkor se szeretnénk a globális memóriából olvasni
• Az első lépés a pixelek beolvasása lokális memóriába – Ez ráadásul két fázisban történik, a kilógó pixelek miatt
• Utána szinkronizáció kell • Végül lehet szűrni 66
GPGPU Szűrés • Szeparábilis szűrők – A szűrő előáll egy horizontális és egy vertikális szűrő szorzataként – SVD-vel lehet meghatározni • Előfordulhat, hogy több szeparábilis szűrő összegeként áll elő – Nem csak egy szinguláris érték van – Ilyenkor nagyság szerint az első néhány értékkel legkisebb négyzetes hiba értelemben lehet legjobb közelítést adni – A szűréseket végezzük két kernel hívással (vertikális és horizontális) és összegezzünk egy harmadik kernellel
– Gauss, Box szűrő ilyen 67
GPGPU Szűrés
http://igm.univ-mlv.fr/~biri/Enseignement/MII2/Donnees/convolutionSeparable.pdf
68
GPGPU Szűrés • • • •
30-as kernelméretnél már nem annyira gyors Szűrjünk frekvenciatérben! FFT is gyorsabb GPU-n FFT + Szorzás + iFFT még mindig megéri – Érdemes az AMD vagy NVidia által adott csomagokat használni
69
GPGPU Naiv Flood fill • Paintes vödör • Lehet tartományt is megadni neki, nem kell teljesen azonos pixelekre folyatni – Watershed innen már nem sok módosítás
• Minden pixlen megvizsgáljuk, hogy ahhoz a színhez tartozik-e amit el akarunk árasztani • Ha igen, akkor megnézzük, hogy van-e elárasztott szomszédja és ha van akkor őt is megjelöljük • Ezt addig ismételjük, amíg van változás a képen 70
GPGPU Flood Fill • Nagyon kevés pixel változik egy lépésben – Rengeteg globális szinkronizáció
• Töltsük be a lokális memóriába a pixeleket – Meg még egy sort a szomszéd blokkokból
• Végezzük el ezen a területen lokális szinkronizációkkal a flood fill-t • Írjuk ki a globális memóriába a saját területet • Ezt ismételgessük – Kb. 200-szor gyorsabb, mint az előző megoldás
• Vannak szofisztikáltabb megoldások a lokális memóriában végzett műveletekre 71
GPGPU Ping-pongozás • Iteratív algoritmusoknál használjuk – A Flood fill is iteratív algoritmus
• • • • •
Érdemes két puffert használni Először az egyikből a másikba dolgozunk Utána a másikból az egyikba dolgozunk Nem kell mindig új puffert foglalni Ha nem 3D textúra, akkor tudunk írható és olvasható puffereket csinálni 72
GPGPU Naiv Parallel reduction • Cél az összes pixel intenzitásának összegét kiszámolni • Adjuk össze párosával a pixeleket és akkor keletkezik egy fele akkora kép amin megismételhetjük a műveletet • Sok globális szinkronizáció
73
GPGPU Parallel reduction • Összegezhetünk a lokális memóriában régiókat • 256-os blokkmérettel 8-ad annyi globális szinkronizáció • Érdemes folytonos blokkokat olvasni a memóriából, az összefésülést kevésbé szereti – Figyeljünk a hatékony indexelésre a globális memória esetén
• Lehet minimumot, maximumot is sok egyebet is számolni ami asszociatív és kommutatív művelet 74
Hisztogram • Minden pixelre megvizsgálom, hogy melyik oszlophoz tartozik • Az oszlopot növelem egyel • Atomi műveletek kellenek! – NVidia Maxwell GPU-k óta (900-as sorozat) van a lokális memóriába is hardware-es atomi művelet
• Érdemes a lokális memóriába összegezni képrészleteket és a végén összeadni a kis hisztogramokat 75
Gyorsító kártya technológiák • OpenCL • OpenMP • Intel fordítója a Xeon Phi kártyákhoz, fejlesztés C/C++ nyelven • FPGA – Verilog/Vhdl – CPU oldalon driver
• Intel+Altera FPGA-k a Xeon sorozatban remélhetőleg hamarosan – Az egyik mag helyett FPGA a QPI buszon 76
Xeon Phi • Eredetileg a GPU-k ellen indult – Nem vált be
• Nagyon drága – Tesla kártyákkal versenyben
• Dinamikus párhuzamosság sokkal jobb támogatása • Hagyományos programozási módszerek kihasználhatósága – OpenMP
• OpenCL-el is programozható – Nem támogatja a textúrázást
77
Xeon Phi – Knights Corner • • • •
16Gb memória 50 darab P54C mag (Pentium 1995-ből) Kiegészítve AVX feldolgozással 50x 8 float művelet = 400 művelet egyszerre – Sokkal rugalmasabb, mint a GPU • Egyszerre többféle művelet
78
Xeon Phi – Knights Landing • Berakható a fő processzor helyére • 72 Airmont Atom mag – 4 szál magonként (HT) – 2 Vektor processzor magonként
• AVX-512 utasítások • TSX – Hardware támogatás tranzakcionális memóriához
• 8-16Gb közeli 3D MCDRAM memória • 384Gb távoli DDR4 memória • 2304 float párhuzamosan – Közben ugyan ennyi nem vektorprocesszoros művelet végrehajtása 79
Egyéb optimalizációs tanácsok • Kerüljük az osztást – Na meg a gyökvonást, nem egész kitevős hatványozást, logaritmust, trigonometriát... – Ha el tudod tárolni a reciprokát és szorozni, akkor úgy csináld
• Kerüljük az int, float, double konverziókat • A float gyorsabb, mint a double – A char meg még gyorsabb 80
Egyéb optimalizációs tanácsok • Bonyolult műveleteknél megérheti táblázatból kiolvasni az eredményt számolás helyett • A fordítónál kapcsolj be mindent releasenél – gcc: -O3 – Visual C++: -Ox
• Várakozó ciklusba mindig rakjunk valami sleep utasítást – Különben nagyon pörgeti a processzort és nem tud mást csinálni 81
Egyéb optimalizációs tanácsok • Használjuk az inline függvényeket – A fordítót is be lehet állítani, hogy automatikusan inlineoljon
• Használjunk a __cdecl hívásó konvenció helyett hatékonyabbat ahol lehet – __fastcall, __vectorcall
• Sebesség szempontjából kritikus helyeken kerüljük a fölösleges objektumokat – Így nem kell konstruktort és destruktort hívni
• Használjunk jól optimalizált könyvtárakat – FFTW, Intel MKL,... 82
Egyéb optimalizációs tanácsok • Írjunk move konstruktort ha szükséges – C++11
• Használjuk a const kulcsszót – Sokat segít a fordítónak
• Kerüljük a managed nyelveket a kritikus részeknél – A Java, C# mindig lasabb lesz, mint a C++ – Lehet modulonként keverni a nyelveket 83
