GPGPU és számítások heterogén rendszereken

GPGPU és számítások heterogén rendszereken

Eichhardt Iván  [email protected]

ELTE-s GPGPU óra anyagai  http://cg.inf.elte.hu/~gpgpu/



Gyors bevezetés a Párhuzamosságról



OpenCL API



Gyakorlati példák (C++)



Párhuzamos programozás elméleti megközelítésben



Lassú a fény?



Kevés a tranzisztor?



Netán rosszul van megírva?



Kezdjük egy kis számolási példával  A fény sebessége 300.000 km/s.  Van egy gépünk egy 3,5 GHz-es CPU-val. ▪ A proci 2 órajel alatt képes lebegőpontos összeadásra

 1 méteres kábellel csatlakoztatunk egy USB

merevlemezt. ▪ A következő adat a lemezen van, be kell olvasni.

 Hány összeadást tudna a proci addig elvégezni, amíg

a lemezről az elektronok (~fénysebességgel) megérkeznek a CPU-ba?



A processzor órajele 3.5Ghz = 3.500.000.000Hz  Egy órajel 1/3.500.000.000 = 285,7 psec (pikosec).  Két órajel (egy összeadás) 571,8 psec időt vesz igénybe.



A fény ennyi idő alatt 3*10^8m/s *571,8 *10^-12s = 0.17m utat tesz meg.

 Méteres kábel esetén: 1m/0,17m ~=~ 6 utasításnyi idő.



Tehát a CPU-nk 6 utasításnyi időt malmozott, amíg az adat megérkezett.



és akkor egyéb lassító tényezővel nem is számoltunk...  Pl. HDD tipikus elérése 20 msec...

...ami alatt a CPU szabadságra is mehet.

 

Az algoritmusaink egymás utáni (szekvenciális) lépésekből állnak. A független lépéseket lehetne egyszerre végrehajtani…  Kulcs: párhuzamosítás.  Ami számtalan módon lehtséges.



Természetesen mindennek megvan az ára.



Lassú a fény?  Utaztassuk kevesebbet az adatokat!  Mozgassunk egyszerre több adatot! (32, 64, …)  Gyorsítótár! (Cache)



Kevés a tranzisztor?  Párhuzamos architektúrák!









SISD – Single Instruction Single Data  Minden utasítás a saját adatával foglalkozik.



MIMD – Multiple Instruction Multiple Data  Egyszerre több utasítás dolgozik a saját adatain. ▪ Több processzormag ▪ Több szál / processzormag



MISD – Multiple Instruction Single Data  Robusztusság…



SIMD – Single Instruction Multiple Data  Ugyan azt az utasítást több adaton végezzük el.  GPU archihektúra…



Csíkszélesség csökkentése.  Több tranzisztor fér el egy helyen.  ~14 nm 
 Korlátok: ▪ Atomi méreteken: szivárgó áram. (atom átmérője ~ 10 – 100 pikométer)



Nagy fejlődésen mentek át.  Nagyon specializált, nem programozható

hardware. …  Általános célú számításokra alkalmazható, programozható hardware.



Nagy fejlődésen mentek át.



GPGPU  General-purpose computing on Graphics

Processing Units 

Egy átlagos felhasználó számára a legnagyobb számítási teljesítmény a GPU-ból nyerhető ki.



Ereje a párhuzamosságban rejlik.



Fehérjék feltekeredésének (folding) szimulációja  Folding@home  H1N1 szimuláció



Az elveszett Apollo 11 videó  Források: Felülírt videó, néhány átvett,

és egy felvétel egy monitorról, amin a videót játsszák le éppen.  100x-os gyorsítás 

Neuronhálók tanítása  Deep Learning, stb.



FPGA, GPGPU, CPU  Field-programmable gate array (FPGA)  DES dekódolás esete „Data Encryption Standard” ▪ CPU: 16 millió kulcs / másodperc ▪ GPU: 250millió kulcs / másodperc (GTX-295) ▪ FPGA: ~1.8 milliárd kulcs / másodperc



H1N1 szimuláció L. Barney - Studying the H1N1 virus using NVIDIA GPUs, Nov 2009. http://blogs.nvidia.com/ntersect/2009/11/studying-the-h1n1-virus-usingnvidia-gpus-.html



Apollo 11 R. Wilson - DSP brings you a high-definition moon walk, Sep 2009. http://www.edn.com/article/CA6685974.html



DES dekódolás Dr. Dobbs - Parallel algorithm leads to crypto breakthrough, Jan 2010. http://www.ddj.com/222600319



A GPGPU problémái

A. Ghuloum - The problem(s) with GPGPU, Oct 2007. http://blogs.intel.com/research/2007/10/the_problem_with_gpgpu.php

  

Nincs szinkronizáció és kommunikáció Csővezeték alkalmazása Párhuzamosítás



Alapműveletek: Map, Amplify, Reduce, Sum

Forrás: http://cg.iit.bme.hu/portal/node/311



SIMD  GPU multiprocesszor

(pl. Vertex tulajdonság streamek)  CPU kiterjesztések (SSE*, 3DNow!, MMX, …)  Adatközpontúság, erőteljesen párhuzamosítható 

Az adatot vektorként kezeljük  Például (vec_res, v1, v2 4×32 bites float vektorok): vec_res.x = v1.x + v2.x; vec_res.y = v1.y + v2.y; vec_res.z = v1.z + v2.z; vec_res.w = v1.w + v2.w;  Egy művelettel írható le



32-bit hosszú bináris sztringek Manhattan távolsága



Ciklussal (Szekvenciális megoldás) int bitcount_naive(int x) { int count = 0; while (x != 0) { if ((x & 1) == 1) { count++; } x >>= 1; } return count; }

 32-bit hosszú bináris sztringek

Manhattan távolsága Bármilyen meglepő, egy processzormagon is tudunk “párhuzamosítani”!



32-bit hosszú bináris sztringek Manhattan távolsága



„Párhuzamos” megoldás unsigned int bitcount(unsigned int x) { x = (x & (0x55555555)) + ((x >> 1) & (0x55555555)); x = (x & (0x33333333)) + ((x >> 2) & (0x33333333)); x = (x & (0x0f0f0f0f)) + ((x >> 4) & (0x0f0f0f0f)); x = (x & (0x00ff00ff)) + ((x >> 8) & (0x00ff00ff)); x = (x & (0x0000ffff)) + ((x >> 16) & (0x0000ffff)); return x; }



128-bit hosszú bináris sztringek Manhattan távolsága unsigned int bitcount_128(unsigned int4 x) { const unsigned int4 a1(0x55555555, 0x55555555, const unsigned int4 a2(0x33333333, 0x33333333, const unsigned int4 a3(0x0f0f0f0f, 0x0f0f0f0f, const unsigned int4 a4(0x00ff00ff, 0x00ff00ff, const unsigned int4 a5(0x0000ffff, 0x0000ffff, x = (x x = (x x = (x x = (x x = (x return }

& (a1)) + & (a2)) + & (a3)) + & (a4)) + & (a5)) + x.x + x.y

((x >> 1) & (a1)); ((x >> 2) & (a2)); ((x >> 4) & (a3)); ((x >> 8) & (a4)); ((x >> 16) & (a5)); + x.z + x.w;

0x55555555, 0x33333333, 0x0f0f0f0f, 0x00ff00ff, 0x0000ffff,

0x55555555); 0x33333333); 0x0f0f0f0f); 0x00ff00ff); 0x0000ffff);



Sok 128 hosszú bit-sztringre: K-NearestNeighbours



1 maggal, naiv megoldással: Lassú 



1 maggal + SIMD: Párhuzamos: sokszoros gyorsulás 



8 maggal CPU-n:



még gyorsabb..

GPU-val (soksok mag): > NAGYON gyors! 

GPGPU 

Csak GPU-k

HETEROGENEOUS COMPUTING 

 Stream programozás  Compute Shader

 CUDA  stb.  (hardverközelibb)



Több mint 1 fajta processzor (CPU, GPU, ...) OpenCL szabvány  Nyílt



Adat- és feladat párhuzamos modell



Az OpenCL nyílt szabvány  A Khronos Group felügyeli



Az OpenCL-C nyelv  ISO C99 szabvány részhalmaza



Numerikus műveletek az IEEE754 alapján



Heterogén platform támogatás  A modell alkalmazható a modern GPU-kra, CPU-kra, Cell

processzorra, DSP-kre, Intel Xenon Phi, Altera FPGA stb...



Mezei felhasználó számára is elérhető!!! (CPU / GPU)      

AMD (OpenCL >v2.0) ARM (OpenCL >v2.0) Intel (OpenCL >v2.0) NVIDIA (OpenCL v1.2) … (Androidon is)

Mi csak OpenCL v1.2-el foglalkozunk.

Mi csak OpenCL v1.2-el foglalkozunk.



Az OpenCL elemei  Platform modell ▪ A Host és az eszköz kapcsolata  Program modell ▪ Data-parallel és Task-parallel lehetőségek  Végrehajtási séma  Memória modell




  

Hoszt eszköz (Host) OpenCL eszköz (Compute Device, „Device”) Számolási egység (Compute Unit, „CU”)  Pl.: NVidia kártyák multiprocesszora



Feldolgozó egység (Processing Element, „PE”)  Pl.: Videokártya Stream processzora  Pl.: CPU magja






Data-parallel és Task-parallel lehetőségek



Data-parallel (Adat párhuzamos) modell  Adat-feladat egység összerendelés  Műveletsor végrehajtása több adatra  A végrehajtás automatikus elosztása



Task-parallel (Feladat párhuzamos) modell  Több független feladat párhuzamos végrehajtása






Host feladata  Kontextus kezelés  Végrehajtás vezérlés



Kernel program  Számító Egységek vezérlése  Egy munkacsoporton belül azonos feladat

elvégzésére



Kernel program  Feladat egységek (Work-Items) ▪ Globális azonosító (global ID) ▪ Minden munkacsoportban azonos program ▪ Egységenként eltérhet a vezérlés  Munkacsoportok (Work-Groups)  Index tér (NDRange)



Kernel program  Feladat egységek (Work-Items)  Munkacsoportok (Work-Groups) ▪ Munkacsoport azonosító (work-group ID) ▪ A feladat egységeknek lokális azonosító (local ID)  Index tér (NDRange)



Kernel program  Munkacsoportok (Work-Groups)  Feladat egységek (Work-Items)  Index tér (NDRange) ▪ N dimenziós problématér N = 1, 2, 3 ▪ Adott méretek N dimenzióban: ▪ Globális címtér ▪ Munkacsoport méret

▪ Azonosítók / indexelés N dimenzióban : ▪ Global ID [pl.: get_global_id(1)] ▪ Local ID [pl.: get_local_id(0)]



Kontextus (Context)  Eszközök (Device)  Kernelek (OpenCL függvények)  Program objektumok (Program) ▪ Forrás ▪ Végrehajtható bináris  Memória objektumok ▪ A Host és az Eszközök által használt memória ▪ A Kernelek ezt látják



Parancs sorok (command-queue)  Host által ellenőrzött  Kernelek végrehajtását vezérli  Parancsok ▪ Kernel végrehajtás ▪ Memória műveletek (írás és olvasás) ▪ Szinkronizáció

 In-order / Out-of-order végrehajtási módok






Négy memória régió az Eszközön  Globális  Konstans  Lokális

 Privát



Globális memória  Írható / olvasható bármelyik Work-Itemből  Bármely eleme elérhető bármely PE-ből  A Host… ▪ foglalja le (allokálja) a területet, ▪ végzi a másolást, ▪ és a memória felszabadítást.



Konstans memória  A Globális memória csak olvasható megfelelője  Kernelben statikusan is definiálható  Néhány hardware külön erre a célra fenntartott,

hatékony memóriaterülettel rendelkezik.



Lokális memória  A Host nem fér hozzá  Egy WG osztott memóriája ▪ Minden WI olvashatja/írhatja



Privát memória  A Host nem fér hozzá

 Csak az adott WI látja



A konzisztenciáról (Hol zavarnak be egymásnak a memória műveletek?)

 Work-Item szinten ▪ WI-ek között nem konzisztens, ▪ De egy WI-en belül konzisztens.  Work-Group szinten ▪ Konzisztens Lokális és Globális memória egy WG-on belül ▪ Nem konzisztens a Globlális memória a WG-ok között



WorkGroup szinkronizáció  WI-ek szinkronizációja  Barrier ▪ Blokkoló hívás

 WG-ok között nincs szinkronizáció!!!! 

CommandQueue szinkronizáció  Parancssori Barrier ▪ garantált a barrier előtti parancsok lefutása ▪ CQ-k között nincs szinkronizáció

 Várakozás Eseményre ▪ minden parancs generál egy eseményt, erre várakoztathatunk egy másikat

Az OpenCL C nyelv

  

C99 nyelv módosítva Skalár típusok Vektor típusok (n ∈ {2,4,8,16})     



(u)charn (u)shortn (u)intn (u)longn floatn

Vektor komponensek elérése (float4 f;)  Swizzle (f.xyzw, f.x, f.xy, f.xxyy, stb.)  Numerikus indexek  Felezés (f.odd, f.even, f.lo, f.hi)



Implicit konverzió  Korlátozott mértékben használható; skalártípusok

között 

Explicit konverzió (Példák)  float4 f = (float4)1.0;  uchar4 u;

int4 c = convert_int4(u);  float f = 1.0f;

uint u = as_uint(f); // 0x3f800000 lesz az értéke



Memóriaterület-jelölők  __global, __local, __constant, __private ▪ Például: __global float4 color;



Függvény-jelölők  __kernel

Egy OpenCL C függvényt Kernelnek jelöl ki.  __attribute__ Fordító segítő attribútumok.



Beépített függvények a végrehajtási modellre vonatkozóan  get_work_dim()  size_t get_{global/local}_{id/size}(uint dimIdx); ▪ Pl.: size_t id = get_global_id(1);  size_t get_num_groups(uint dimIdx);  size_t get_group_id(uint dimIdx);



Szinkronizációs függvények  barrier(flag); ▪ CLK_LOCAL_MEM_FENCE : lokális memóriát konzisztensé teszi ▪ CLK_GLOBAL_MEM_FENCE : globális memóriát konzisztensé teszi  mem_fence(flag);  write_mem_fence(flag);  read_mem_fence(flag);



További beépített függvények  Általános, szokásos függvények  Geometriai függvények  Összehasonlító függvények floatn típusokon ▪ (isequal, isinfinite, stb)  Memóriára vonatkozóan: ▪ Aszinkron memória olvasás ▪ Prefetch (cache-be töltés globális memóriából)



VexCL  C++ template könyvtár vektor-kifejezések írására és futtatására

OpenCL/CUDA-n  https://github.com/ddemidov/vexcl 

SharpCL  C# bytecoderól OpenCL-re fordít és futtat  Érdeklődni az előadónál.



HadoopCL  http://pubs.cs.rice.edu/node/336  MapReduce Heterogén rendszereken, Hadoop és OpenCL

integrációval



Apple OS X Snow Leopard  http://en.wikipedia.org/wiki/Mac_OS_X_Snow_Leopard#OpenCL



…



Esettanulmányok: Programozási minták (Elmélettel)  Map  (Gather)

 (Scatter)  Stencil  Reduce

 Scan



„Leképezés”  Többváltozós függvény alkalmazása.



Általában a többi tervezési mintával együtt használják.  Így új tervminták jöhetnek létre.



Az implementációról  Helyben is végezhető (a bemeneten).

P

P

P

Gather: A „P” feladat több helyről tetszése szerint gyűjt adatokat, melyből egy elemet hoz létre a kimeneten. P

P

Scatter: A „P” folyamat tetszése szerint több kimeneti elemet érint.

P

Például: A scatter során a „P” folyamat egy elem (módosított) értékét több másikhoz adja.



~„Sablon”  Adott N hosszú input és N hosszú output.  Az éppen feldolgozott elem valamilyen környezete alapján

számolunk outputot.  1D, 2D, 3D, …. 

Alkalmazás  Pl.: Futószűrő (helyben végzett)  Pl.: Elmosás (nem helyben végzett) ▪ Doboz szűrő, Gauss szűrő, …



Technikák  Konvolúciók, Medián, Véges differenciák,

Bilaterális szűrés, stb…



Implementáció  Konvolúcióknál: Szeparábilis? => gyorsítás



Fontos szempont:  A stencil egy fix minta alapján gyűjti az adatokat.  A stencil minden kimenetre ír.



Naiv megvalósításának problémája:  Adat-újrafelhasználás.  ▪ Kérdés: Mi a megoldás pl. átlagszűrő esetén?

 Táblás feladat: ▪ Kérdés: Hányszor fogunk olvasni egy-egy input adatelemet az adott minta alapján?



Kérdés: Milyen tervezési minta alapján működhet a következő…  Összegzés?

 Feltételes értékadás?  Rendezés?



Map:



 Index tér 

 Index tér  Bemenet

(Bemenet, Kimenet)  „One-to-One”

 „One-to-Many”

 

Scatter:

Stencil:

Gather:

 Index tér 

 Index tér  Kimenet

(Bemenet*, Kimenet)  „Several-to-One”

 „Many-to-One”

~”Tömörítés”  Adott művelet: ⊕ ▪ Asszociatív. ▪ Kommutatív?  Adott bemenet, ahol értelmezett a művelet.  A bemenetet a művelet segítségével redukáljuk

egyetlen egységgé: ▪ Pl.: Számsorozat összege (művelet: összadás) ▪ Pl.: Maximumkiválasztás

Alkalmazások  Összegzés, Maximum-kiválasztás …  Csak-asszociatív esetben: Scan tervezési minta

egy lépése! („Up-Sweep”) 

MapReduce  http://research.google.com/archive/mapreduce.html  Nagy klasztereken nagyon klassz!

Implementáció  Végezhető helyben, vagy „váltogatott” bemenet-

kimenet tömbökön.  Szinkronizáció szükséges a hatékony megvalósításhoz.  Ha a műveletre teljesül:

▪ Asszociativitás: Alapeset. ▪ Kommutativitás: Bizonyos hardvereken előnyös lehet, ha a kommutativitást kihasználva egymás mellé rakjuk a lépésenként redukált sorozatokat. ▪ Nem maradnak „közök” az elemek között. ▪ SIMD

Értékek: Lépés 1

„get_id”

0

1

2

3

4

5

6


„get_id” 0

1

2


„get_id”

0

Értékek:

Lépés 4

„get_id” 0 Értékek:

1

3

7


„get_id”

0

1

2

3

0

1

2

3

0

1


„get_id” Értékek:

Lépés 3


Lépés 4


0

4

5

6

7

„Asszociatív művelet eredménye” (fogalom: ORSI-tantárgy) 

Adott művelet: ⊕  Asszociatív.



A művelet ismételt alkalmazásánál a részeredményeket is kiszámítjuk.  Összehasonlítva: A „Reduce” leginkább csak a

végeredménnyel törődik.

Alkalmazások   

A Radix rendezés egyik lépése. Változó szélességű képszűrés. Adatfolyam feldolgozás.

Irodalom:  Blelloch, Guy E. 1990. "Prefix Sums and Their Applications." Technical Report CMU-CS-90-190, School of Computer Science, Carnegie Mellon University.  http://people.inf.elte.hu/hz/parh/parhprg.html

Implementáció  Naiv megvalósítás: Nem „munka”-hatékony: O(n log2 n)  Hatékony megvalósítás: ▪ „Kiegyensúlyozott fák” mintájára (balanced trees) ▪ PP-ben hasznos algoritmikus minta!!! ▪ Bináris fa.

▪ A bemenet hosszával azonos nagyságrendű hatékonyság: O(n). ▪ A bináris fát nem tároljuk le, csak az elvet használjuk! ▪ Két lépésben: ▪ „Up-Sweep” fázis (Reduce minta) ▪ „Down-Sweep” fázis

▪ Irodalom: Belloch (1990)


„get_id”

0

1

2

3

4

5

6

0

1

2

3

4

5

0

1

2

3



Lépés 3




Az algoritmus két lépésben működik:  Up-Sweep  Down-Sweep



Up-Sweep  Egy csak-asszociatív „Reduce”.



Down-Sweep Az utolsó elem kinullázása. A redukciós lépések végrehajtása fordított sorrendben: Nagy lépésektől a kisebbekig. Egy lépésre:

1. 2. 1. 2. 3.



Eredmény := művelet(Bal oldali elem, Jobb oldali elem) (Redukció) Bal oldali elem := Jobb oldali elem. (Ez egy „új” lépés) Jobb oldali elem := Eredmény. (Redukció befejezése)

Más néven: Blelloch Scan

Köszönöm a figyelmet!

GPGPU és számítások heterogén rendszereken

Recommend Documents