CUDA: více o sdílené paměti

Rekapitulace

Sdružený přístup do paměti II.

Ukázali jsme si, jak řešit sdružený přístup do paměti, pokud naše proměnná nemá velikost 4, 8, 16 bytů. Jako příklad byla použita práce s float3 polem. Pole typu float3 jsme při volání kernelu převedli na pole float uložené do sdílené paměti. Poté se ze sdílené paměti “vytahovala” float3 čísla, se kterými jsme pracovali dle zadání a zpět uložili do sdílené paměti jako typ float. Závěrem se předalo pole float ze sdílené paměti na výstup. U struktur porušující velikost 4, 8, 16 bytů použijeme zarovnání pomocí klíčového slova __align(X)__, kde X je požadovaná velikost struktury, tedy 4, 8 nebo 16. Rovněž jsme si v souvislosti se sdruženým přístupem do paměti ukázali, že je lepší používat strukturu polí než pole struktur, a to z důvodu snadnější manipulace s daty.

Debuggování CUDA programů

K debuggování se ladí pomocí debuggeru cuda-gdb CUDA program, který je velmi podobný klasickému gdb (doplněn o několik příkazů vhodných pro CUDA). Při práci s cuda-gdb nesmí grafické zařízení, které bude použito pro úlohu debuggování, používat zároveň Xka. Je tomu tak z důvodu ochrany, kdy Xka mohou při špatně napsaném CUDA programu “zatuhnout”. Pokud máme jedno GPU, musí být Xka před debuggováním vypnuta. V opačném případě nám cuda-gdb nebude sloužit.

Debuggování CUDA programů II.

Kromě dobře známých gdb příkazů: ®un, (b)reak, ©ontinue, (s)top, (n)ext, (p)rint, (d)elete, (i)nfo (nebudou zde rozebírána – pokud čtenář nezná základy gdb, jistě na internetu najde mnoho kvalitních tutoriálů), můžeme použít v CUDA programech thread, info cuda threads a info cuda state.

1. Thread:

   Příkaz slouží k libovolnému přepínání bloků a vláken. Má několik podob:

thread <<<(BX,BY),(TX, TY, TZ)>>>

(mezera mezi slovem thread a <<< je nutná!)

thread <<<(X),(Y)>>>

thread <<<(X)>>>

thread

Thread <<<(BX,BY),(TX, TY, TZ)>>> je   úplný zápis, kde BX, BY definujeme souřadnice bloku a TX, TY a TZ definujeme
   souřadnice vlákna BX, BY bloku. Po příkazu thread můžeme vytisknout příkazem print    libovolnou proměnnou, jejíž hodnota bude odpovídat požadovanému vláknu.

V případě varianty thread<<<(X),(Y)>>> bude X x-ová složka bloku a Y bude x-ová složka vlákna. Ostatní složky jsou nastaveny na 0 (obdobně je tomu tak při volání kernelu)! Varianta thread<<<(X)>>> (může být i bez závorek) definujeme pouze x-ovou složku vlákna. Blok v tomto případě zůstane stejný. Poslední variantou thread bez parametrů vypíšeme aktuální blok a vlákno.

2. info cuda threads:

   vypíše informace o počtu bloků a vláken ve tvaru:
   <<<(BX0,BY0),(TX0,TY0,TZ0)>>> … <<<(BX1,BY1),(TX1, TY1, TZ1)>>>, kde složky  končící 0 reprezentují počáteční blok a vlákno, složky končící 1 reprezentují poslední blok a  vlákno. Jinými slovy, vypíše rozsah bloků a vláken na základě definovaného gridDim a blockDim v programu (při volání kernelů). Akceptovatelná varianta tohoto příkazu je info cuda threads all.

3. info cuda state:

Prostřednictvím tohoto příkazu se vypíší informace o použitém hardwaru a rovněž i o paměti alokované pomocí funkce cudaMalloc.

Nyní malá ukázka práce s cuda-gdb (použijeme zdrojový kód nocoalesced.cu z minulého dílu):

1. Nejdříve zkompilujeme program pomocí nvcc s parametry -g -G:

   nvcc -g -G noncoalesced.cu -o noncoalesced

2. Spustíme cuda-gdb (připomínám, že pokud máme jen jedno GPU musí být Xka vypnuta):

   cuda-gdb noncoalesced

3. Nastavíme break na začátek kernelu:

   b noncoalesced_float3_kernel

4. Spustíme program:

   r

5. Vypíšeme si built-in proměnné blockDim, blockIdx, threadIdx

   p blockDim
   Výpis:
   $1 = { x = 64, y = 1, z = 1}

   p blockIdx
   Výpis:
   $2 = { x = 0, y = 0}

   p threadIdx
   Výpis:
   $3 = { x = 0, y = 0, z = 0}

6. Vypíšeme aktuální vlákno a blok:

   thread
   Výpis:
     [Current Thread 2 (Thread -1210931504 (LWP 25528))]
    [Current CUDA Thread <<<(0,0),(0,0,0)>>>]

7. “Skočíme” na řádek 10
    b 10
    Výstup:
     Breakpoint 2 at 0x8064050: file noncoalesced.cu, line 10

8. Vypíšeme proměnnou idx:
    p idx
    Výstup:
     $4 = 0

9. Přepneme vlákno na souřadnice (3, 0, 0) v aktuálním bloku (ještě jednou upozorňuji na  mezeru po slově thread):

    thread <<<(3)>>>

10. Znovu vypíšeme proměnnou idx:

    p idx
    Výpis:
    $5 = 3

Vidíme, že se již nevypisuje hodnota vlákna (0,0,0), ale vlákna (3,0,0) aktuálního bloku (0,0).

11. Na závěr vypíšeme rozsah vláken a bloků a informace o alokované paměti (pomocí cudaMalloc) a o použité verzi architektury:

    info cuda threads
    Výstup:
    <<<(0,0),(0,0,0)>>>...<<<(7,0),(63,0,0)>>> noncoalesced_float3_kernel() at
    noncoalesced.cu: 6
    info cuda state

    (moc dlouhý výstup :-)

Více o sdílené paměti

O sdílené paměti víme, že je velmi rychlá (v oficiální dokumentaci [1] se uvádí rychlost srovnatelná s registry), jak ji alokovat pro naše programy a že je dobrým pomocníkem při sdruženém přístupu do globální paměti (pokud není dodržena velikost paměťového elementu na 4, 8 nebo 16 bytech). Co je důležité vědět o sdílené paměti? Se sdílenou pamětí je provázána existence bank a konfliktů
bank.

Banky jsou malé paměťové elementy (přístup do banky zabere 2 cykly) o velikosti 32bitů (4 byty).Nedojde-li ke konfliktu bank, lze k nim přistupovat současně prostřednictvím half-warpu. V rámci celého warpu se jedná o 2 “transakce”, kde jedna je pro dolní a druhá pro horní half-warp. Konflikt bank spočívá v tom, že dvě paměťové adresy v rámci half-warpu odkazují do stejné banky. V takovém případě se přístup do bank serializuje a není již současný. Konfliktu bank se říká n-pásmové bankovní konflikty (n-way bank conflicts), kde n označuje počet paměťových adres přistupujících do jedné banky.

Existuje jedna vyjímka, kdy nedojde ke konfliktu bank, a to když přistupuje všech 16 vláken half-warpu ke stejné bance.

Nyní si ukažme několik možných situací a také, jak se vyvarovat případným konfliktům bank:

1. Nejpoužívanějším případem je přístup do pole (např. float nebo int), kde index se rovná id   vlákna:

   __shared__ float shared[32];
   float data = shared[BaseIndex + s * tid];

   V tomto případě nedochází k žádnému konfliktu, protože každé vlákno přistupuje do jiné banky (obr. 5.1).

2. Dalším případem může být přístup k paměťovým elementům menším než jsou 4 byty.  Například u pole typu char:

   __shared__ char shared[32];
   char data = shared[BaseIndex + tid];

Tento přístup je konfliktní. Pro eliminaci konfliktu bank můžeme přistupovat k jednotlivým  elementům jako 4-násobek id vlákna:

   char data = shared[BaseIndex + 4 * tid];

3. Posledním používaným případem je přístup k paměťovým elementům větších jak 4 byty, například double (8 bytů). Abychom se vyhnuli konfliktu bank (konkrétně 2-way bank conflict) rozdělíme proměnnou typu double na 2 integery (tedy 2× 4 byty). K tomu použijeme built-in funkce v kernelu jako __double2loint, __double2hiint a __hiloint2double  takto:

__shared__ int shared_lo[32];
__shared__ int shared_hi[32];
double dataIn;
shared_lo[BaseIndex + tid] = __double2loint(dataIn);
shared_hi[BaseIndex + tid] = __double2hiint(dataIn);
double dataOut =
    __hiloint2double(shared_hi[BaseIndex + tid],
                      shared_lo[BaseIndex + tid]);

4. Obdobně je to se strukturami, kde musíme dávat pozor na jejich velikost.

Na závěr si ukážeme grafické znázornění přístupu do bank s/bez konfliktů.

Obr. 5.1: Přístup do bank sdílené paměti bez konfliktu bank

Obr. 5.2: Přístup do bank sdílené paměti s konfliktem bank

První část obrázku (vlevo) ukazuje 2-way bank conflicts. Na druhé části obrázku (vpravo) je zachycen 8-way bank conflicts.

Závěr

V příštím, již závěrečném díle o základech architektury CUDA, si představíme profileru, který může být stejně tak dobrým pomocníkem jako cuda-gdb a také několik dalších rad a tipů pro vývoj CUDA programů.

Použitá literatura:
[1] NVIDIA_CUDA_Pro­gramming_Guide_2­.2.1.pdf
[2] nvcc₂.2.pdf
[3] CUDA_GDB_User_Ma­nual.pdf