Maximizing GPU Performance with Atomic Operations in CUDA Programming

1. CUDA C/C++ programozás Atomikus műveletek A segédanyag készítése a TÁMOP 4.2.4.A/2-11-1-2012-0001 Nemzeti Kiválóság Program című kiemelt projekt keretében zajlott. A projekt az Európai Unió támogatásával, az Európai Szociális Alap társfinanszírozásával valósul meg.

2. GPU számítási képességek • Számítási képesség ismétlés: • Minden GPU-hoz tartozik egy adott számítási képesság. • Pl.: 1.2, 2.0, 3.0 • Leírja: • Az architektúra főbb jellemzőit. • A GPU által adott lehetőségeket, és szolgáltatásokat. • Inkrementális szolgáltatáslista. • A magasabb számítási képességű GPU több szolgáltatást ad. • Régi elemek nem vesznek el. • Visszafelé kompatibilitás.

3. Fordítás adott számítási képességre • A kód fordításakor jelezni lehet, a fordító számára, hogy: • Mi a minimális szolgáltatáskészlet (számítási képesség) amire fordítani szeretnénk a programot. • Milyen GPU-ra optimalizáljon. • Megadás az nvcc parancssori argumentumával: • nvcc–arch=sm_11 • A program minimum 1.1-es képességű GPU-n fut. • Miért érdekes ez nekünk? • Mert a mai óra anyaga, az atomikus műveletek a 2.0-s képességtől jönnek be a GPU-ba. • Meg különben sem árt tudni.

4. Konkurens futtatás • A párhuzamos programok futásakor előfordulhatnak versenyhelyzetek: • Két vagy több utasítás ugyanazt az adatot próbálja elérni/módosítani. • Például: __global__ void kernel(int* x) { (*x)++; // x növelése (nem pixelenként) return; } Int main(...) { int* dev_x; // ... kernel<<<2, 2>>>(dev_x); // melyik szál növeli x-et? // ... }

5. Probléma konkurens adatelérésnél • (*x)++; • 3 lépésben • X értékének kiolvasása a memóriából, • Az érték növelése, • Eredmény visszaírása a memóriába • Ha párhuzamosan megy (pl.: x): • 1. szál • Olvasás: • reg = *x;// 4 beolvasva • Növelés: • reg++;// növelt érték: 5 • Visszaírás: • *x = reg; // *x=5; • 2. szál • Olvasás: • reg = *x;// 4 beolvasva • Növelés: • reg++;// növelt érték: 5 • Visszaírás: • *x = reg; // *x=5; • 3. szál • Olvasás: • reg = *x;// 4 beolvasva • Növelés: • reg++;// növelt érték: 5 • Visszaírás: • *x = reg; // *x=5; idő

6. Ugyanaz szekvenciálisan • (*x)++; • 3 lépésben • X értékének kiolvasása a memóriából, • Az érték növelése, • Eredmény visszaírása a memóriába • Ha párhuzamosan megy (pl.: x): • 1. futás • Olvasás: • reg = *x;// 4 beolvasva • Növelés: • reg++;// növelt érték: 5 • Visszaírás: • *x = reg; // *x=5; • 2. futás • Olvasás: • reg = *x;// 5 beolvasva • Növelés: • reg++;// növelt érték: 6 • Visszaírás: • *x = reg; // *x=6; • 3. futás • Olvasás: • reg = *x;// 6 beolvasva • Növelés: • reg++;// növelt érték: 7 • Visszaírás: • *x = reg; // *x=7; idő

7. Konkurens változónövelés a gyakorlatban • 14_RaceCondition.cu • Egy változó érték növelése adott számszor. • Párhuzamosan a GPU-n. • És szekvenciálisan a GPU-n.

8. Atomikus műveletek • Feloldják a konfliktusokat. • Biztosított, hogy egyszerre csak egy szál férhessen hozzá egy adatelemhez. • Műveletek: • int atomicAdd(int* address, intval); • intatomicSub(int* address, intval); • int atomicExch(int* address, intval); • int atomicMin(int* address, intval); • int atomicMax(int* address, intval); • unsigned intatomicInc(unsigned int* address,unsigned intval); • unsigned intatomicDec(unsigned int* address,unsigned intval); • intatomicCAS(int* address, int compare, intval); • intatomicAnd(int* address, intval); • int atomicOr(int* address, intval); • int atomicXor(int* address, intval);

9. Valós példa:Hisztogram • Hisztogram számítása a GPU-n. • Adott egy adatsor (diszkrét elemekből). • Pl.: kép egész pixelintenzitásokkal 0 és 255 között. • Megoldása CPU-n. • 15_HistCPU.cu • Sima ügy.  unsignedchar* adat; int* hiszt[256]; for(int i=0; i<256; i++) hiszt[i] = 0; for(int i=0; i<adat_meret; i++) hiszt[adat[i]]++;

10. Hisztogram a GPU-n • Bonyolult művelet. • A szálak konkurens módon párhuzamosan dolgozzák fel az adatot. • Konfliktus léphet fel, ha két szál ugyanazt a hisztogram számlálót akarja növelni. • Magoldás: atomikus műveletek • A számláló növeléseket atomikus műveletekkel végezzük a kernelben. • Ez automatikusan feloldja a konfliktusokat. • Hátrányok: • Lassítja a kódot • A szálak sorban várhatnak egy számláló növelésére. • 16_Hist_GPU.cu

11. További lehetőségek • Mire lehetne még használni az atomikus műveleteket?