11. fejezet: GPU-kutatási irányok a szkalarizáció és az affin végrehajtás területén
Ahogy a 2. fejezetben is leírtuk, a GPU-számítási API-k, mint a CUDA és az OpenCL, egy MIMD-szerű programozási modellt kínálnak, amely lehetővé teszi a programozó számára, hogy nagy tömb méretű skalár szálakat indítson a GPU-ra. Bár minden egyes skalár szál követheti a saját egyedi végrehajtási útvonalát, és hozzáférhet tetszőleges memóriahelyekhez, a leggyakoribb esetben mindannyian egy kis végrehajtási útvonal-készletet követnek, és hasonló műveleteket végeznek.
A GPU-szálak konvergens vezérlési folyama a legtöbb, ha nem az összes modern GPU-n ki van használva a SIMT végrehajtási modell révén, ahol a skalár szálakat warp-okba csoportosítják, amelyek SIMD hardveren futnak (lásd a 3.1.1. szakaszt). Ez a fejezet egy kutatási sorozatot foglal össze, amely tovább kihasználja ezeknek a skalár szálaknak a hasonlóságát a szkalarizáció és az affin végrehajtás révén.
Ennek a kutatásnak a kulcsfontosságú felismerése a [Kim et al., 2013] által leírt értékstruktúra megfigyelésén alapul, amely a ugyanazon számítási kernel végrehajtása során lévő szálak között figyelhető meg. Az értékstruktúra két típusa, az egyenletes és az affin, a 11.1. példában bemutatott számítási kernelben szemléltetett.
Egyenletes változó
Olyan változó, amely ugyanazt az állandó értéket veszi fel minden szál esetében a számítási kernelben. A 11.1. algoritmusban a a változó, valamint a THRESHOLD és Y_MAX_VALUE literálok mind egyenletes értékekkel rendelkeznek az összes szál esetében a számítási kernelben. Egy egyenletes változó egyetlen skalár regiszterben tárolható, és minden szál által újrafelhasználható a számítási kernelben.
Affin változó
Olyan változó, amelynek az értékei a szál azonosítójának lineáris függvényei minden szál esetében a számítási kernelben. A 11.1. algoritmusban a y[idx] változó memóriacíme az idx szál azonosítójának affin transzformációjaként ábrázolható:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Ez az affin reprezentáció két skalár értékként, egy bázisként és egy lépésközként tárolható, ami sokkal kompaktabb, mint a teljesen kibontott vektor.
__global__ void vsadd( int y[], int a ) {
// Kód
}
```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.
nt idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
11.1 algoritmus: Példa skalár és affin műveletek egy számítási kernelben ([Kim et al., 2013] alapján).
Több kutatási javaslat is van arra, hogyan lehet észlelni és kihasználni az egyenletes vagy affin változókat a GPU-kon. A fejezet hátralévő része ezeket a javaslatokat foglalja össze ezekben a két aspektusban.
## Egyenletes vagy affin változók észlelése
A GPU számítási kernelben lévő egyenletes vagy affin változók létezésének észlelésére két fő megközelítés van: Fordító-vezérelt észlelés és Hardveres észlelés.
### Fordító-vezérelt észlelés
Egy lehetséges mód a GPU számítási kernelben lévő egyenletes vagy affin változók létezésének észlelésére egy speciális fordítói elemzés végrehajtása. Ez lehetséges, mert a meglévő GPU programozási modellek, a CUDA és az OpenCL, már biztosítanak módot a programozó számára, hogy egy változót állandónak nyilvánítson a számítási kernel teljes futása során, valamint egy speciális változót a szál azonosítójához. A fordító végezhet egy vezérlési függőségi elemzést, hogy észlelje azokat a változókat, amelyek kizárólag állandóktól és szál azonosítóktól függenek, és megjelölje őket egyenletesként/affinként. Az egyenletes/affin változókon kizárólag működő műveletek ezután jelöltek a szkalarizálásra.
Az AMD GCN [AMD, 2012] a fordítóra támaszkodik az egyenletes változók és a dedikált skalár processzor által tárolható és feldolgozható skalár műveletek észlelésére.
Asanovic et al. [2013] egy kombinált konvergens és változó elemzést vezetnek be, amely lehetővé teszi a fordító számára, hogy meghatározza egy tetszőleges számítási kernelben azokat a műveleteket, amelyek alkalmasak szkalarizálásra és/vagy affin transzformációra. A számítási kernel konvergens régióin belüli utasítások skalár/affin utasításokká alakíthatók. Bármely divergens és konvergens régiók közötti átmenetkor a fordító beszúr egy `syncwarp` utasítást a vezérlési folyamat által indukált regiszterfüggőségek kezelésére. Asanovic et al. [2013] átvették ezt a megközelítést.Itt a magyar fordítás a megadott markdown fájlhoz. A kódban nem fordítottam le a kódot, csak a megjegyzéseket.
his analysis to generate scalar operations for the Temporal-SIMT architecture [Keckler et al., 2011, Krashinsky, 2011].
Ez az elemzés a Temporal-SIMT architektúrához [Keckler et al., 2011, Krashinsky, 2011] tartozó skalár műveletek generálására szolgál.
Decoupled Affine Computation (DAC) [Wang and Lin, 2017] relies on a similar compiler analysis to extract scalar and affine candidates to be decoupled into a separate warp. Wang and Lin [2017] augment the process with a divergent affine analysis, with the goal to extract strands of instructions
A Decoupled Affine Computation (DAC) [Wang and Lin, 2017] hasonló fordítói elemzésre támaszkodik, hogy kinyerje a skalár és affin jelölteket, amelyeket külön warphozrendelnek. Wang és Lin [2017] a folyamatot divergens affin elemzéssel egészítették ki, azzal a céllal, hogy utasítássorozatokat nyerjenek ki.