7. fejezet: Streaming Multiprocessor tervezés GPU-tervezésben

A streaming multiprocessor (SM) a NVIDIA GPU-architektúrák alapvető építőeleme. Mindegyik SM egy sor CUDA-mag-ot tartalmaz, amelyek utasításokat SIMT (Single Instruction, Multiple Thread) módon hajtanak végre. Az SM felel a warpokok kezeléséért és ütemezéséért, az ágazati divergencia kezeléséért és a megosztott memória és gyorsítótárak gyors eléréséért. Ebben a fejezetben megvizsgáljuk az SM mikroarchitektúráját, beleértve a csővezetékeket, a warp ütemezési mechanizmusokat, a regiszterfájl-tervezést, valamint a megosztott memória és az L1 gyorsítótár szervezését.

SM mikroarchitektúra és csővezetékek

Az SM egy erősen párhuzamos és csővezetékes processzor, amely arra tervezték, hogy hatékonyan hajtson végre több száz szálat egyidejűleg. A 7.1. ábra a NVIDIA Volta-architektúra egy egyszerűsített blokkvázlatát mutatja.

                                 Utasítás-gyorsítótár
                                         |
                                         v
                                    Warp-ütemező
                                         |
                                         v
                               Diszpécser egység (4 warp)
                                 |   |   |   |
                                 v   v   v   v
                               CUDA-mag (FP64/FP32/INT)
                               CUDA-mag (FP64/FP32/INT)
                               CUDA-mag (FP64/FP32/INT)
                               ...
                               Tensor-mag
                               Tensor-mag
                               ...
                               Terhelés/Tárolás egység
                               Terhelés/Tárolás egység
                               ...
                               Speciális függvény egység
                                         ^
                                         |
                                Regiszterfájl (64 KB)
                                         ^
```Hungarian translation:

Megosztott memória / L1 cache (96 KB)

7.1. ábra: Az NVIDIA Volta architektúra SM-jének egyszerűsített blokkdiagramja.

Az SM fő komponensei közé tartoznak:

1. Utasítás-gyorsítótár: Gyakran elért utasításokat tárol a késleltetés csökkentése és a teljesítmény javítása érdekében.

2. Warp-ütemező: Kiválasztja a végrehajtásra kész warp-okat, és elküldi őket a rendelkezésre álló végrehajtó egységekhez.

3. Végrehajtó egység: Beolvassa és dekódolja a legfeljebb 4 warp/ciklus utasításait, majd elküldi őket a megfelelő végrehajtó egységekhez.

4. CUDA-magok: Programozható végrehajtó egységek, melyek széles körű egész- és lebegőpontos műveleteket támogatnak. Minden Volta SM-ben 64 CUDA-mag található.

5. Tensor-magok: Speciális végrehajtó egységek, melyek a mélytanulási és AI-terhelések gyorsítására lettek tervezve. Minden Volta SM-ben 8 Tensor-mag található.

6. Betöltő/Tároló egységek: Kezelik a memória-műveleteket, beleértve a globális memória, a megosztott memória és a gyorsítótárak be- és kiírását.

7. Speciális funkciós egységek: Transzcendens és egyéb összetett matematikai műveleteket hajtanak végre.

8. Regiszterfájl: Gyors hozzáférést biztosít a szál-privát regiszterekhez. Minden Volta SM-ben 64 KB-os regiszterfájl található.

9. Megosztott memória / L1 gyorsítótár: Egy konfigurálható memóriatér, amely szoftver által kezelt gyorsítótárként (megosztott memória) vagy hardver által kezelt L1 adatgyorsítótárként használható.

Az SM-pipeline úgy van tervezve, hogy maximalizálja a teljesítményt, lehetővé téve több warp egyidejű végrehajtását és a memórialatencia elrejtését. A 7.2. ábra az SM-pipeline egy egyszerűsített nézetét mutatja.

    Utasítás-beolvasás
            |
            v
    Utasítás-dekódolás
            |
            v
    Operandus-gyűjtés
            |
            v
    Végrehajtás (CUDA-magok, Tensor-magok, Betöltő/Tároló egységek, Speciális funkciós egységek)
            |
            v
    Visszaírás

7.2. ábra: Az SM-pipeline egyszerűsített ábrája.

A pipeline-szakaszok a következők:

1. Utasítás-beolvasás: A warp-ütemező kiválasztja a végrehajtásra kész warp-ot, ésKérem, itt van a magyar fordítás a megadott markdown fájlhoz. A kódban nem fordítottam le a kódot, csak a kommenteket.

2. Utasítás-előfetch: Az SM előfetch-eli a következő utasítást az aktuális warphoz az utasítás-gyorsítótárból.

3. Utasítás-dekódolás: A letöltött utasítás dekódolásra kerül, hogy meghatározzuk a művelettípust, a bemeneti operandusokat és a célregisztereket.

4. Operandus-gyűjtés: A szükséges operandusok összegyűjtésre kerülnek a regiszterfájlból vagy a megosztott memóriából.

5. Végrehajtás: Az utasítás végrehajtatik a megfelelő végrehajtási egységen (CUDA mag, Tensor mag, Terhelés/Tárolás egység vagy Speciális Függvény egység).

6. Visszaírás: A végrehajtás eredménye visszaírásra kerül a regiszterfájlba vagy a megosztott memóriába.

A nagy teljesítmény elérése érdekében az SM több technikát is alkalmaz az erőforrás-kihasználtság maximalizálására és a késleltetés elrejtésére:

• Kettős kiadás: Az SM két független utasítást is kiadhat egy warpnak egyetlen ciklusban, lehetővé téve a magasabb szintű utasításszintű párhuzamosságot.
• Csővezérelt végrehajtási egységek: A végrehajtási egységek csővezéreltek, lehetővé téve az SM számára, hogy új műveletet kezdjen egy egységen, még mielőtt az előző művelet befejeződött volna.
• Késleltetés elrejtése: Az SM ciklusonként válthat a warpok között, lehetővé téve számára, hogy elrejtse a memória-hozzáférések és a hosszú késleltetésű műveletek késleltetését más warpok utasításainak végrehajtásával.

A 7.1. példa egy egyszerű CUDA kernelt mutat be, amely két vektor elemenként! összegzését végzi.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

7.1. példa: CUDA kernel vektor-összeadáshoz.

Ebben a példában minden szál a kernelben kiszámítja a bemeneti vektorok a és b megfelelő elemeinek összegét, és az eredményt eltárolja a c kimeneti vektorban. Az SM ezt a kernelt úgy hajtja végre, hogy minden szálat egy CUDA maghoz rendel, és warpokat ütemez a rendelkezésre álló magok végrehajtására. A terhelés/tárolás egységeket használják a bemeneti adatok globális memóriából való betöltésére és az eredmények visszaírására.

Warp-ütemezés és divergencia-kezelés

EfHatékony warp ütemezés kulcsfontosságú a SM teljesítményének maximalizálása szempontjából. A warp ütemező felel a végrehajtásra kész warps kiválasztásáért és az elérhető végrehajtó egységekre való elküldéséért. A warp ütemező elsődleges célja, hogy a végrehajtó egységeket elfoglalva tartsa, biztosítva, hogy mindig legyenek végrehajtásra kész warps.

A SM kétszintű warp ütemezési mechanizmust alkalmaz:

1. Warp ütemezés: A warp ütemező kiválasztja a végrehajtásra kész warpokat egy ütemezési politika, például körforgás vagy legöregebb először alapján. A kiválasztott warpokat ezután elküldik a rendelkezésre álló végrehajtó egységekre.

2. Utasítás ütemezés: Az egyes warpok esetén a SM az utasítások függőségei és a végrehajtó egységek rendelkezésre állása alapján ütemezi az utasításokat. A SM képes több független utasítás egyidejű kiadására ugyanabból a warpból, hogy maximalizálja az utasítás szintű párhuzamosságot.

A 7.3. ábra a kétszintű warp ütemezési mechanizmust szemlélteti.

    Warp készlet
    Warp 1 (Kész)
    Warp 2 (Várakozó)
    Warp 3 (Kész)
    ...
    Warp N (Kész)
        |
        v
    Warp ütemező
        |
        v
    Kiadó egység
        |
        v
    Végrehajtó egységek

7.3. ábra: Kétszintű warp ütemezési mechanizmus.

A warp ütemezés egyik kulcskihívása a ágeltérés kezelése. A SIMT végrehajtási modellben a warp minden szála egy időben hajtja végre ugyanazt az utasítást. Azonban, amikor egy warp elágazás utasításba (pl. if-else) ütközik, egyes szálak a if-ágat, mások pedig az else-ágat követik. Ezt a helyzetet ágeltérésnek nevezzük.

Az ágeltérés kezelésére a SM egy technikát alkalmaz, amely predikáció néven ismert. Amikor egy warp eltérő ágba ütközik, a SM mindkét ágat egymás után hajtja végre, elrejtve azokat a szálakat, amelyek nem követik az adott ágat. Az eredményeket azután predikátumregiszterek segítségével egyesítik, hogy minden szál a helyes eredményt kapja.

A 7.2. példa egy CUDA kernel kódját mutatja be ágeltéréssel.Itt a magyar fordítás a megadott Markdown fájlhoz. A kódban lévő megjegyzéseket fordítottam le, a kódot nem változtattam meg.

__global__ void divergentKernel(int *data, int *result) {
    // A threadId kiszámítása a blokk indexe és a thread indexe alapján
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) {
        // Ha a data[tid] nagyobb, mint 0, akkor szorozzuk meg 2-vel
        result[tid] = data[tid] * 2;
    } else {
        // Egyébként szorozzuk meg 3-mal
        result[tid] = data[tid] * 3;
    }
}

7.2. példa: Divergens ágat tartalmazó CUDA kernel.

Ebben a példában az data[tid] > 0 elágazási feltétel miatt előfordulhat, hogy egy warp egyes szálai a if ágat, míg mások az else ágat hajtják végre. Az SM ezt a divergenciát úgy kezeli, hogy mindkét ágat egymás után, szekvenciálisan hajtja végre, és a nem aktív szálakat kikapcsolja.

A 7.4. ábra szemlélteti a predikáció folyamatát egy divergens szálakkal rendelkező warp esetén.

    Warp (32 szál)
    1. szál: data[1] = 5, result[1] = 10
    2. szál: data[2] = -3, result[2] = -9
    ...
    32. szál: data[32] = 7, result[32] = 14

    Divergens elágazás:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Predikáció:
    1. lépés: Hajtsd végre az `if` ágat a maszk alapján
        1. szál: result[1] = 10
        2. szál: (inaktív)
        ...
        32. szál: result[32] = 14

    2. lépés: Hajtsd végre az `else` ágat a maszk alapján
        1. szál: (inaktív)
        2. szál: result[2] = -9
        ...
        32. szál: (inaktív)

    Végleges eredmény:
    1. szál: result[1] = 10
    2. szál: result[2] = -9
    ...
    32. szál: result[32] = 14

7.4. ábra: Predikáció folyamata egy divergens szálakkal rendelkező warp esetén.

A predikáció használatával az SM kezelheti a divergens elágazásokat explicit ági utasítások vagy vezérlési művelet-divergencia nélkül. Azonban a divergens elágazások még mindig hatással lehetnek a teljesítményre, mivel az SM mindkét ágat egymás után, szekvenciálisan hajtja végre, ami csökkenti a tényleges párhuzamosságot.

Regiszterfájl és operandus-gyűjtők

A regiszterfájl a SM kulcsfontosságú komponense, amely gyors hozzáférést biztosít a szál-privát regiszterekhez. Minden SM-nek nagy regiszterfájlja van, hogy támogassa a sok aktív szálat, és lehetővé tegye a warptok hatékony kontextusváltását.Itt az a magyar fordítás a fájlra, a kódban csak a megjegyzéseket fordítottam le:

A NVIDIA Volta architektúrában minden SM-nek (streaming multiprocessor) 64 KB-os regiszterfájlja van, amely 32 db 2 KB-os bankból áll. A regiszterfájl úgy van kialakítva, hogy nagy sávszélességet és alacsony késleltetést biztosítson a nagy számú egyidejű szál támogatására.

A bankütközések minimalizálása és a teljesítmény javítása érdekében az SM egy úgynevezett operandus-gyűjtő (operand collector) technikát alkalmaz. Az operandus-gyűjtők speciális egységek, amelyek összegyűjtik az operandusokat a regiszterfájl bankjaiból és továbbítják őket a végrehajtó egységekhez. Az operandus-gyűjtők használatával az SM csökkentheti a bankütközések hatását és javíthatja a végrehajtó egységek kihasználtságát.

A 7.5. ábra egy egyszerűsített diagramot mutat be az SM regiszterfájljáról és operandus-gyűjtőiről.

    Regiszterfájl (64 KB)
    Bank 1 (2 KB)
    Bank 2 (2 KB)
    ...
    Bank 32 (2 KB)
        |
        v
    Operandus-gyűjtők
        |
        v
    Végrehajtó egységek

7.5. ábra: Regiszterfájl és operandus-gyűjtők egy SM-ben.

Az operandus-gyűjtők több utasítás és több warp operandusait gyűjtik össze, lehetővé téve az SM számára, hogy egyetlen ciklusban különböző warpokból származó utasításokat bocsásson ki a végrehajtó egységekre. Ez segít elrejteni a regiszterfájl-hozzáférések késleltetését, és javítja az SM általános átviteli teljesítményét.

A 7.3. példa egy CUDA kernelt mutat be, amely két vektor skaláris szorzatát számítja ki.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // A szálak részleges összegeket számítanak
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    // A részleges összegek összegzése
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    // Az összeg kiírása
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

Ebben a példában minden szál a saját részleges összegét számítja ki a hozzá rendelt elemekre. Ezután a részleges összegek összefűzése és összegzése következik.Here is the Hungarian translation of the provided file, with the code comments translated:

elemek a bemeneti vektorokból. A részleges összegek a partialSum megosztott memória tömbben vannak tárolva. Miután minden szál kiszámította a saját részleges összegeit, párhuzamos csökkentést hajtanak végre a részleges összegek összegzésére, hogy megkapják a végleges skaláris szorzat eredményt.

Az operandus-gyűjtő kulcsfontosságú szerepet játszik ebben a példában, mivel hatékonyan gyűjti össze az operandusokat a megosztott memória-hozzáférésekhez és aritmetikai műveletekhez. Segít elkerülni a bank-konfliktusokat, és javítja a végrehajtási egységek kihasználtságát.

Következtetés

A streamingmultiprocesszor a modern GPU-architektúrák fő számítási egysége. Tervezése a propellersebesség maximalizálására és a memória-késleltetés elrejtésére összpontosít a finom szemcsés multithreading, a SIMT-végrehajtás és a hatékony operandus-gyűjtés kombinációjával.

Az SM kulcsfontosságú összetevői közé tartozik a warp-ütemező, amely kiválasztja a végrehajtandó warpokat; a SIMT-verem, amely kezeli az ágak divergenciáját és konvergenciáját; a regiszterállomány és az operandus-gyűjtők, amelyek gyors hozzáférést biztosítanak a szálfüggő regiszterekhez; valamint a megosztott memória és az L1 cache, amelyek alacsony késleltetésű adatmegosztást és újrafelhasználást tesznek lehetővé.

Ahogy a GPU-architektúrák tovább fejlődnek, a kutatás olyan területeken, mint az ágdivergencia-kezelés, a warp-ütemezés és a regiszterállomány tervezése, kulcsfontosságú lesz a jövőbeli GPU-k teljesítményének és hatékonyságának javítása érdekében. Az olyan újszerű technikák, mint a dinamikus warp-képzés, a threadblokk-tömörítés és az operandus-újrafelhasználási gyorsítótárak, jelentősen növelhetik az SM képességeit, és új szintű teljesítményt biztosíthatnak a párhuzamos számítási terhelésekben.