2. fejezet: GPU programozási modellek
A Grafikus Feldolgozó Egységek (GPU-k) fejlődése során a fix funkcionalitású grafikus gyorsítóktól eljutottak a nagyon párhuzamos, programozható számítási motorokig, melyek képesek a különböző alkalmazások széles körének gyorsítására. A GPU-k masszív párhuzamosságának hatékony kiaknázása érdekében számos párhuzamos programozási modell és API fejlődött ki, mint például az NVIDIA CUDA, az OpenCL és a DirectCompute. Ezek a programozási modellek absztrakciós rétegeket biztosítanak, amelyek lehetővé teszik a programozók számára, hogy kifejezhessék az alkalmazásaikban rejlő párhuzamosságot, miközben elrejtik a GPU hardver alacsony szintű részleteit.
Ebben a fejezetben megismerkedünk a GPU-k számára készült párhuzamos programozási modellek mögötti kulcsfontosságú koncepciókkal és elvekkel, a végrehajtási modellre, a GPU utasításkészlet architektúrákra (ISA-k), az NVIDIA GPU ISA-ra és az AMD Graphics Core Next (GCN) ISA-ra összpontosítva. Példákat is bemutatunk, amelyek illusztrálják, hogyan alkalmazzák ezeket a koncepciókat a gyakorlatban.
Végrehajtási modell
A modern GPU programozási modellek végrehajtási modellje a kernelekre épül, amelyek olyan függvények, amelyeket párhuzamosan hajt végre a GPU-n lévő nagyszámú szál. Egy kernel indításakor a programozó megadja a létrehozandó szálak számát és azok hierarchikus elrendezését rácsokba, tömbökbe (vagy kooperatív szálcsoportokba - CTA) és egyedi szálakba.
- Egy rács képviseli a teljes problémateret és egy vagy több tömbből áll.
- Egy tömb a szálak egy csoportja, amelyek együttműködhetnek és szinkronizálhatják magukat megosztott memória és korlátok segítségével. A tömbön belüli szálakat ugyanazon a GPU magban (streaming multiprocessor vagy compute unit) hajtják végre.
A szálaknak egyedi azonosítójuk van a tömbön és a rácsön belül, amely felhasználható a memóriacímek kiszámításához és az irányítási folyamat döntéseihez.
Ez a hierarchikus szervezés lehetővé teszi a programozók számára, hogy mind az adatszintű párhuzamosságot (amikor ugyanazt a műveletet több adatelemre alkalmazzák), mind a feladatszintű párhuzamosságot (amikor különböző feladatok hajtódnak végre párhuzamosan) kifejezhessék.
ÁbraKérem, itt van a fájl magyar fordítása. A kódnál nem fordítottam le a kódot, csak a megjegyzéseket.
Rács
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Blokk |
| | | |
Szál Szál ...2.1 ábra: Szálhierarchia a GPU végrehajtási modellben.
SIMT végrehajtás
A CUDA és OpenCL GPU programozási modellek Egyetlen Utasítás, Több Szál (SIMT) végrehajtási modellt követnek. A SIMT modellben a szálak csoportokban, ún. warp-okban (NVIDIA terminológia) vagy wavefront-okban (AMD terminológia) hajtódnak végre. A warp minden szála ugyanazt az utasítást hajtja végre egyszerre, de mindegyik szál más adaton dolgozik.
Azonban, eltérően a hagyományos Egyetlen Utasítás, Több Adat (SIMD) modeltől, ahol minden feldolgozási elem zárt sorban működik, a SIMT lehetővé teszi, hogy a szálak független végrehajtási útvonalakkal rendelkezzenek és elágazásoknál eltérő utakat válasszanak. Amikor egy warp elágazási utasításba ütközik, a GPU hardver kiértékeli az elágazási feltételt minden warp-ban lévő szálra. Ha minden szál ugyanazt az utat választja (konvergált), a warp normálisan folytatja a végrehajtást. Ha néhány szál más utat választ (divergált), a warp két vagy több alwarp-ra oszlik, mindegyik egy másik útvonalat követve. A GPU hardver szekvenciálisan hajtja végre a divergens utakat, kikapcsolva a nem aktív szálakat mindegyik alwarp-ban. Amikor minden út befejeződik, az alwarp-ok újra konvergálnak és zárt sorban folytatják a végrehajtást.
A 2.2 ábra bemutatja a SIMT végrehajtást divergens vezérlési folyammal.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Elágazás |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
Újrakonvergencia2.2 ábra: SIMT végrehajtás divergens vezérlési folyammal.
Ez a divergencia kezelési mechanizmus lehetővé teszi, hogy a SIMT rugalmasabban támogassa a vezérlési folyamatokat, mint a hagyományos SIMD modell.Kérjük, itt van a fájl magyar fordítása. A kódnál ne fordítsd le a kódot, csak a megjegyzéseket fordítsd le.
HAN SIMD, de ez a SIMD hatékonyság csökkenésének költségével jár, amikor divergencia történik. A programozóknak törekedniük kell a divergencia minimalizálására egy warp-on belül az optimális teljesítmény elérése érdekében.
Memória-hierarchia
A GPU-k egy összetett memória-hierarchiával rendelkeznek, hogy támogassák a párhuzamos munkafeladatok magas sávszélességű és alacsony késleltetésű követelményeit. A memória-hierarchia általában a következőket tartalmazza:
- Globális memória: A legnagyobb, de leglassúbb memóriatér, amely minden szál számára elérhető egy kernelben. A globális memória általában nagy sávszélességű GDDR vagy HBM memóriával van megvalósítva.
- Megosztott memória: Egy gyors, lapkára integrált memóriatér, amelyet egy blokk összes szála megoszt. A megosztott memóriát a szálak közötti kommunikációra és az adatok megosztására használják egy blokkon belül.
- Állandó memória: Egy csak olvasható memóriatér, amelyet az összes szál számára sugárzott, csak olvasható adatok tárolására használnak.
- Textúra memória: Egy csak olvasható memóriatér, amely a térbeli lokalitásra van optimalizálva, és textúra gyorsítótárakon keresztül érhető el. A textúra memóriát általában grafikai munkafeladatokban használják.
- Helyi memória: Egy privát memóriatér minden szál számára, amely a regiszter kihelyezésére és a nagy adatszerkezetekre szolgál. A helyi memória általában a globális memóriára van leképezve.
A memória-hierarchia hatékony kihasználása kulcsfontosságú a GPU-k magas teljesítményének eléréséhez. A programozóknak arra kell törekedniük, hogy maximalizálják a megosztott memória használatát, és minimalizálják a globális memóriához való hozzáféréseket a memória-késleltetés és sávszélesség-szűk keresztmetszetek csökkentése érdekében.
A 2.3. ábra a GPU memória-hierarchiáját mutatja.
____________
| |
| Globális |
| Memória |
____________
|
____________
| |
| Állandó |
| Memória |
____________
|
____________
| |
| Textúra |
| Memória |
____________
|
|
____________
| |
| Megosztott |
| Memória |
____________
|
____________
| |
| Helyi |
| Memória |
____________ÁbraKérjük, adja meg a következő markdown fájl magyar fordítását. A kód esetében ne fordítsa le a kódot, csak a megjegyzéseket fordítsa le. Itt van a fájl: ure 2.3: GPU memória hierarchia.
GPU utasításkészlet-architektúrák
A GPU utasításkészlet-architektúrák (ISA-k) meghatározzák a szoftver és a hardver közötti alacsony szintű interfészt. Meghatározzák a GPU által támogatott utasításokat, regisztereket és memória-címzési módokat. A GPU ISA-k megértése létfontosságú a hatékony GPU-kód fejlesztéséhez és a teljesítmény optimalizálásához.
Ebben a részben két vezető GPU-gyártó, az NVIDIA és az AMD ISA-it fogjuk feltárni. Az NVIDIA Parallel Thread Execution (PTX) és SASS ISA-ira, valamint az AMD Graphics Core Next (GCN) ISA-jára fogunk összpontosítani.
NVIDIA GPU ISA-k
Az NVIDIA GPU-k két szintű ISA-t támogatnak: PTX (Parallel Thread Execution) és SASS (Streaming ASSembler). A PTX egy virtuális ISA, amely stabil célpontot biztosít a CUDA fordítók számára, míg a SASS az NVIDIA GPU-k natív ISA-ja.
PTX (Parallel Thread Execution)
A PTX egy alacsony szintű, virtuális ISA, amelyet az NVIDIA GPU-khoz terveztek. Hasonló az LLVM IR-hez vagy a Java byte-kódhoz abban, hogy egy stabil, architektúra-független célpontot biztosít a fordítók számára. A CUDA programokat általában PTX kódba fordítják, amelyet aztán az NVIDIA GPU-illesztőprogram natív SASS utasításokká alakít.
A PTX széles körű aritmetikai, memória- és vezérlési utasításokat támogat. Korlátlan számú virtuális regisztert és predikációt támogat, ami lehetővé teszi a vezérlési folyamatok hatékony végrehajtását. A PTX speciális utasításokat is biztosít a szál-szinkronizációhoz, az atomos műveletekhez és a textúra mintavételezéshez.
Itt egy példa egy egyszerű vektor-összeadó kernel PTX kódjára:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
cvta.to.global.u64 %rd4, %rd1;
cvta
```Itt látható a magyar fordítás a kódhoz:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Magyar fordítás a kommentárokhoz:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
// Globális memóriából olvasás az f1 és f2 lebegőpontos változókba ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; // Az f1 és f2 értékek összeadása, az eredmény az f3 lebegőpontos változóba add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; // A globális memóriába írás az f3 értékkel st.global.f32 [%rd10], %f3;
// Visszatérés a függvényből ret; }
A kód egy vektorösszeadó kernel függvényt definiál, amely négy paramétert kap: a bemeneti és kimeneti vektorok címeit, valamint a vektorok méretét. A kernel kiszámítja a globális szálon belüli azonosítót, betölti a megfelelő elemeket a bemeneti vektorokból, végrehajtja az összeadást, és eltárolja az eredményt a kimeneti vektorban.
A SASS (Streaming ASSembler) a NVIDIA GPU-k natív ISA-ja, amely közvetlenül leképeződik a GPU hardverre. A SASS-utasítások a NVIDIA GPU-illesztőprogram által generálódnak a PTX-kódból, és általában nem láthatók a programozók számára.
Az AMD Graphics Core Next (GCN) architektúra és ISA a AMD GPU-kat használja. A GCN egy RISC-alapú ISA, amely támogatja mind a grafikai, mind a számítási feladatokat. Magas teljesítményt, skálázhatóságot és energiahatékonyságot céloz meg.Itt a magyar fordítás a megadott markdown fájlhoz. A kódban csak a megjegyzéseket fordítottam le, a kódot nem.
alar ALU hatékony skaláris műveletek és vezérlési áramlás végrehajtásához.
- Vektor ALU párhuzamos, adatpárhuzamos műveletek végrehajtásához.
- Nagy sávszélességű memóriarendszer, amely támogatja az atomi műveleteket és az alacsony késleltetésű hozzáférést a megosztott memóriához.
- Rugalmas címzési mód a memóriaműveletekhez, amely támogatja az alap+eltolás és a skaláris+vektor címzést.
Itt egy példa a GCN ISA kódra egy vektor-összeadási kernel esetén:
```asm
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
s_load_dwordx4 s[0:3], s[4:5], 0x0 # Argumentumok betöltése a memóriából
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
v_lshlrev_b32 v0, 2, v0 # Indexek eltolása
v_add_u32 v1, vcc, s1, v0 # Első vektor címének kiszámítása
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc # Második vektor címének kiszámítása
flat_load_dword v1, v[1:2] # Első vektor betöltése
v_add_u32 v3, vcc, s0, v0 # Harmadik vektor címének kiszámítása
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc # Negyedik vektor címének kiszámítása
flat_load_dword v0, v[3:4] # Második vektor betöltése
v_add_f32 v0, v0, v1 # Vektor-összeadás
flat_store_dword v[3:4], v0 # Eredmény tárolása
s_endpgmEz a GCN kód betölti a bemeneti vektorelemeket flat_load_dword használatával, végrehajtja az összeadást v_add_f32 segítségével, és eltárolja az eredményt a memóriában flat_store_dword használatával. Az s_load_dwordx4 és s_load_dword utasításokat a kernel argumentumok memóriából való betöltésére használják.
Algoritmusok leképezése GPU architektúrákra
Az algoritmusok hatékony leképezése a GPU architektúrára kulcsfontosságú a magas teljesítmény eléréséhez. Fontos szempontok:
Elegendő párhuzamosság biztosítása
Az algoritmust sok, finom szemcsés szálra kell bontani, hogy azok párhuzamosan tudjanak futni, és kihasználják a GPU párhuzamos feldolgozási képességeit. Ez gyakran azt jelenti, hogy meg kell találni az algoritmus adatpárhuzamos részeit, amelyek függetlenül, különböző adatelemeken hajthatók végre.
Elágazás-divergencia minimalizálása
A warp/wavefront-on belüli divergens vezérlési áramlás szerializációhoz és csökkent SIMD-hatékonysághoz vezethet. Az algoritmusokat úgy kell kialakítani, hogy minimalizálják az elágazás-divergenciát, amennyire lehetséges. Ezt úgy lehet elérni, hogy csökkentjük az adattól függő elágazások használatát### A memória-hierarchia kiaknázása
A globális memória elérése költséges. Az algoritmusoknak maximalizálniuk kell a megosztott memória és a regiszterek használatát, hogy csökkentsék a globális memória-hozzáféréseket. Az adatokat is úgy kell elrendezni a memóriában, hogy lehetővé tegyék a koaleszkált memória-hozzáféréseket, ahol a warp-on belüli szálak összefüggő memóriahelyekhez férnek hozzá. A memória-hierarchia hatékony használata jelentősen csökkentheti a memória-késleltetést és a sávszélesség-szűk keresztmetszeteket.
A számítás és a memória-hozzáférések kiegyensúlyozása
Az algoritmusoknak magas arányt kell mutatniuk az aritmetikai műveletek és a memória-műveletek között, hogy hatékonyan elrejthessék a memória-késleltetést, és magas számítási teljesítményt érjenek el. Ezt úgy lehet elérni, hogy maximalizáljuk az adatok újrafelhasználását, előre beolvassuk az adatokat, és átfedésbe hozzuk a számítást a memória-hozzáférésekkel.
A gazdagép-eszköz adatátvitel minimalizálása
Az adatok gazdagép (CPU) és eszköz (GPU) memóriája közötti átvitele lassú. Az algoritmusoknak minimalizálniuk kell ezeket az átviteleket, és a lehető legtöbb számítást a GPU-n kell végezniük. Az adatokat nagy adagokban kell átvinni a GPU-ra, és a lehető legtovább ott kell tartani, hogy kompenzálják az átvitel költségeit.
A GPU-kerneleket fejlesztésénél általánosan használt párhuzamos algoritmus-tervezési minták:
-
Leképezés: Minden szál ugyanazt a műveletet végzi egy különböző adatelemen, lehetővé téve a nagy adathalmazok egyszerű párhuzamos feldolgozását.
-
Csökkentés: A párhuzamos csökkentést arra használják, hogy hatékonyan kiszámítsanak egyetlen értéket (pl. összeg, maximum) egy nagy bemeneti adathalmazból. A szálak helyi csökkentéseket végeznek, amelyeket aztán egyesítenek a végső eredmény előállításához.
-
Scan: Más néven előtagösszeg, a scan arra szolgál, hogy kiszámolja az elemek futó összegét egy tömbben. A hatékony párhuzamos scan-algoritmusok kulcsfontosságú építőelemek sok GPU-gyorsított alkalmazáshoz.
-
Sablonszerűség: Minden szál egy érték kiszámítását végzi a szomszédos adatelemek alapján. A sablonszerű számítások gyakoriak a tudományos szimulációkban és a képfeldolgozási alkalmazásokban.Itt a fájl magyar fordítása:
-
Összegyűjtés/Szétszórás: A szálak tetszőleges helyekről olvasnak (összegyűjtés) vagy írnak (szétszórás) a globális memóriába. A hatékonyság érdekében gondos adatelrendezésre és hozzáférési mintákra van szükség.
A 3.20. ábra egy példát mutat a map módszerre, ahol minden szál egy függvényt (pl. négyzetgyök) alkalmaz a bemeneti tömb egy különböző elemére.
Bemeneti tömb:
| | | | | | | |
v v v v v v v v
______________________________
Szálak: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Kimeneti tömb: 3.20. ábra: A map minta példája GPU programozásban.
Következtetés
A CUDA és az OpenCL, mint GPU programozási modellek, hozzáférést biztosítanak a modern GPU-k párhuzamos feldolgozási képességeihez a fejlesztők számára, lehetővé téve széles körű alkalmazások gyorsítását. Ezek a programozási modellek olyan absztrakciót nyújtanak, amely lehetővé teszi a finoman szemcsézett párhuzamos feladatok hatékony leképezését a GPU hardverre.
A végrehajtási modell, a memóriahierarchia és a szinkronizációs primitívák megértése elengedhetetlen a nagy teljesítményű GPU-kód írásához. A fejlesztőknek gondosan kell figyelembe venniük olyan tényezőket, mint a szálak szervezése, az elágazás divergenciája, a memória-hozzáférési minták és az algoritmus tervezése, hogy teljes mértékben kihasználhassák a GPU-k számításiteljesítmény-képességeit.
Ahogy a GPU-architektúrák tovább fejlődnek, a programozási modellek és eszközök is fejlődniük kell, hogy lehetővé tegyék a fejlesztők számára, hogy hatékonyan hasznosítsák az új hardveres funkciókat és képességeket. A programozási nyelvtervezés, a fordítóprogram-optimalizálás és az automatikus hangolás terén folyó kutatások elengedhetetlenek lesznek a programozói termelékenység és a teljesítményhordozhatóság javításához a heterogén számítástechnika korszakában.