Hlavní navigace

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

Lukáš Zaorálek 18. 8. 2009

Náš seriál o programování paralelních výpočtů probíhajících na grafických kartách se pomalu blíží ke konci. Sdruženou paměť jsme minulým dílem dokončili a dnes nás čeká předposlední díl. Ten naváže na započaté základy debuggování CUDA aplikací a společně v něm probereme sdílenou paměť.

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.cuminulé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ů.

CUDA 5

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

CUDA 5

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.1.pdf
[2] nvcc2.2.pdf
[3] CUDA_GDB_User_Ma­nual.pdf

Našli jste v článku chybu?

18. 8. 2009 22:08

sdfs (neregistrovaný)

zde se potvrzuje, ze po osme hodine vecerni se nema mejlovat ani prispivat do diskuzi. clovek pak lituje co v opilosti napsal :-)

Podnikatel.cz: Přehledná titulka, průvodci, responzivita

Přehledná titulka, průvodci, responzivita

Lupa.cz: Kdo pochopí vtip, může jít do ČT vyvíjet weby

Kdo pochopí vtip, může jít do ČT vyvíjet weby

Podnikatel.cz: Babiše přesvědčila 89letá podnikatelka?!

Babiše přesvědčila 89letá podnikatelka?!

Lupa.cz: Levný tarif pro Brno nebude. Radní: je to kartel

Levný tarif pro Brno nebude. Radní: je to kartel

DigiZone.cz: Recenze Westworld: zavraždit a...

Recenze Westworld: zavraždit a...

Podnikatel.cz: 1. den EET? Problémy s pokladnami

1. den EET? Problémy s pokladnami

Vitalia.cz: Spor o mortadelu: podle Lidlu falšovaná nebyla

Spor o mortadelu: podle Lidlu falšovaná nebyla

DigiZone.cz: Sony KD-55XD8005 s Android 6.0

Sony KD-55XD8005 s Android 6.0

120na80.cz: Pánové, pečujte o svoje přirození a prostatu

Pánové, pečujte o svoje přirození a prostatu

Lupa.cz: Teletext je „internetem hipsterů“

Teletext je „internetem hipsterů“

120na80.cz: Jak oddálit Alzheimera?

Jak oddálit Alzheimera?

Měšec.cz: Golfové pojištění: kde si jej můžete sjednat?

Golfové pojištění: kde si jej můžete sjednat?

Měšec.cz: Air Bank zruší TOP3 garanci a zdražuje kurzy

Air Bank zruší TOP3 garanci a zdražuje kurzy

120na80.cz: Horní cesty dýchací. Zkuste fytofarmaka

Horní cesty dýchací. Zkuste fytofarmaka

Měšec.cz: Jak vymáhat výživné zadarmo?

Jak vymáhat výživné zadarmo?

Vitalia.cz: Říká amoleta - a myslí palačinka

Říká amoleta - a myslí palačinka

Měšec.cz: Jak levně odeslat balík přímo z domu?

Jak levně odeslat balík přímo z domu?

Měšec.cz: U levneELEKTRO.cz už reklamaci nevyřídíte

U levneELEKTRO.cz už reklamaci nevyřídíte

Root.cz: Vypadl Google a rozbilo se toho hodně

Vypadl Google a rozbilo se toho hodně

Lupa.cz: Proč firmy málo chrání data? Chovají se logicky

Proč firmy málo chrání data? Chovají se logicky