Hlavní navigace

CUDA: optimalizace přístupu do globální paměti

4. 8. 2009
Doba čtení: 12 minut

Sdílet

V minulém díle jsme probrali základy CUDA a napsali minimalistický program “Hello CUDA!”. Dnes zrekapitulujeme základy z předešlých dvou dílů a postoupíme dále. Vysvětlíme, k čemu jsou dobré bloky a gridy a ukážeme si optimalizaci přístupu do globální paměti (tzv. sdruženého přístupu do paměti).

Rekapitulace


V předchozích dvou dílech padlo mnoho z terminologie CUDA. Pokusím se zrekapitulovat to
nejdůležitější:

1. Základní architektura CUDA
GPU (označované jako device – zařízení) je tvořeno Multiprocesory.
Multiprocesor (dále jen MP) se zákládá z 8 procesorů.

Paměť GPU je rozdělena následovně:
Každý MP má svou sdílenou paměť (kešovaná), paměť pro textury (kešovaná) a paměť pro konstanty (kešovaná).

Každý procesor obsahuje registry a může přistupovat do paměti (sdílené, textur, konstant) MP, do kterého patří.

Globální paměť (nekešovaná), kam má přistup každý procesor (je společná pro všechny procesory GPU) slouží, jak pro kopírování dat z/do hosta (“CPU”), resp. zařízení(GPU), tak i pro ukládání objemných dat, která se “nevejdou” do sdílené paměti. V takovém případě mluvíme o lokální paměti.

2. Program CUDA:
Kód spouštěný na GPU se nazývá kernel a je reprezentován funkcí jazyka C, která má vždy výstupní hodnotu typu void a identifikátor __global__. CUDA programy je také možno psát v jazyku Fortran. Ten je oficiálně podporovaným jazykem pro CUDA.

Kernel je spouštěn ve vláknech (threads). Vlákna se spouští v tzv. blocích (blocks) a několik bloků tvoří mřížku (grid). Volání kernelu vypadá takto: jmeno_kernelu<<<po­cet_bloku, pocet_vlaken_uv­nitr_bloku>>> (vstupni_para­metry_kernelu);. Pro úplnost dodám, že pocet_bloku
= mřížka a pocet_vlaken_uv­nitr_bloku = jeden blok.

Mřížka, bloky, vlákna


Nyní je na čase si vysvětlit, proč architektura CUDA zdánlivě složitě rozděluje vlákna do bloků (blocks) a mřížek (grid). Důvodem je, že každá série grafických zařízení obsahuje různý počet MP.
Novější série mají samozřejmě větší počet MP než ty starší. Pokud budeme konkrétní, tak G80 série má 16 MP, tedy 128 procesorů (16×8) a 10tková série má již 30 MP, tedy 240 procesorů (30×8). K tomu je třeba říci, že konkrétní zařízení se mohou v počtu MP mírně lišit od své série. To znamená, že musíme najít nějaký efektivní způsob, jak bez nutnosti rekompilace CUDA programu efektivně
využít možnosti staršího, ale i novějšího “železa”. K tomuto účelu slouží mřížka.

Mřížka udává, jak bylo vysvětleno minule, počet bloků. Vlákna uvnitř bloku přistupují do společné sdílené paměti a rovněž mohou být v rámci bloku synchronizována, a to pomocí volání (built-in) funkce __syncthreads() uvnitř kernelu na místech, kde si přejeme synchronizaci vláken. Nutno však upozornit, že např. dvě vlákna, kde 1. vlákno patří do jednoho bloku a 2. vlákno do druhého bloku, nelze synchronizovat! Kromě sdílené paměti mohou vlákna využívat i registry. K registrům lze přistoupit pouze vlákno, nikoliv blok. Pro přehlednost uvádím tabulku 3.1 přístupu k GPU paměti z pohledu kernelu a hosta:

Druh paměti Přístup Umístění Operace Kešovaná
Registry Jedno vlákno Na čipu čtení/zápis Ne
Lokální Jedno vlákno Mimo čip (DRAM) čtení/zápis Ne
Sdílená Všechna vlákna uvnitř jednoho bloku Na čipu čtení/zápis
Globální Všechna vlákna a host Mimo čip (DRAM) čtení/zápis Ne
Pro textury Všechna vlákna a host Mimo čip (DRAM) čtení/zápis Ano
Pro konstanty Všechna vlákna a host Mimo čip (DRAM) čtení/zápis Ano

Tab. 3.1 Rozdělení GPU paměti

Víme, že paměť GPU je rozdělena na několik druhů. Nejdůležitější z nich jsou globální paměť a sdílená paměť.

Globální paměti


Pojďme se podívat na náš druhý program, který bude mít za úkol pracovat s velkým polem čísel (v globální paměti), kde každé číslo vydělíme dvěma.

#include <stdio.h>
__global__ void numbers (int* out, int* in) {
        // spocitame index do pole in a out pro konkretni vlakno
        // pripominam, ze pracujeme pouze s jednorozmernym polem,
        // tudiz jsme definovali x-slozku velikosti mrizky a bloku v execution configuration.
        int idx = blockDim.x * blockIdx.x + threadIdx.x;
        // delime cislo dvema a ukladame do vystupniho pole
        out[idx] = in[idx]/2;
}
int main (int argc, char** argv) {
       //pomocna promenna
       bool result = true;
       //ukazatel na pole integeru pro praci s cisly v kernelu
       int* num_h;
       //ukazatel na pole int v globalni pameti
       int* num_d;
       //ukazatel na vystupni pole v globalni pameti
       int* num_out_d;
       //pozadovana velikost pole
       size_t num_size = 128*512;
       //pocet vlaken na jeden blok
       int num_threads_per_block = 128;
       //velikost mrizky
       int num_blocks = num_size/num_threads_per_block;
       //pozadovana velikost pole v bytech
       size_t num_size_bytes = sizeof (int)*num_size;
       //alokace pameti pole
       num_h = (int*)malloc (num_size_bytes);
       //alokujeme pole num_d v globalni pameti
       cudaMalloc ((void**) &num_d, num_size_bytes);
       //alokujeme vystupni pole num_out_d v globalni pameti
       cudaMalloc ((void**) &num_out_d, num_size_bytes);
       //naplnime pole cisly
       for (unsigned int i = 0; i < num_size; i++) {
          num_h[i] = i*2;
  }
  //kopirovani pole int z hosta do globalni pameti
  cudaMemcpy (num_d, num_h, num_size_bytes, cudaMemcpyHostToDevice);
  //volame kernel
  numbers<<<num_blocks, num_threads_per_block>>> (num_out_d, num_d);
  //cekej, dokud vsechny vlakna neskonci
  cudaThreadSynchronize();
  //nyni muzeme zkopirovat vystupni data num_out_d do num_h
  cudaMemcpy (num_h, num_out_d, num_size_bytes, cudaMemcpyDeviceToHost);
  //zkontrolujeme, zda vysledek je spravny
  for (unsigned int i = 0; i < num_size; i++) {
          if (i != num_h[i]) {
                   result = false;
                   printf ("Vysledek je spatny! %d\n", i);
                   break;
          }
  }
  //pokud je vysledek spravny, vytiskneme oznameni
  if (result) {
          printf ("Vysledek je spravny!\n");
  }
  //uvolneni prostredku globalni pameti
  cudaFree (num_d);
  cudaFree (num_out_d);
  //uvolneni prostredku v pameti hosta
  free (num_h);
  return 0;
}

Kód 3.1 Program numbers.cu

Pracujeme-li s globální pamětí, budeme používat API funkce CUDA jako cudaMalloc, cudaMemcpy a cudaFree (jejich mnohem více, ale prozatím si bohatě vystačíme s těmito základními). Jejich popis uvádím (kvůli odlišnostem od standardních funkcí malloc, memcpy, apod.) v tabulce 3.2.

Jméno funkce Vstupní parametry Výstupní Popis
cudaMalloc void devPtr, size_t size cudaError_t Alokace paměti ukazatele devPtr o velikosti size
cudaMemcpy void dst, const void src, size_t count, enum cudaMemcpyKind kind cudaError_t Kopírování paměti z src do dst o velikosti count podle kind
cudaMemset void devPtr, int value, size_t count cudaError_t Nastavení paměti ukazatele devPtr o velikosti count na hodnotu value
cudaFree void devPtr cudaError_t Uvolnění paměti ukazatele devPtr

Tab. 3.2 Funkce pro práce s globální pamětí

Všechny funkce mají jako výstupní hodnotu cudaError_t, podle které se dá zjistit, zda nedošlo k chybě a pokud ano, k jaké (o chybách později). Funkce cudaMemcpy je trošku rozdílná od klasické funkce memcpy. Z jejího předpisu (cudaError_t cudaMemcpy (void dst, const void src, size_t count,
enum cudaMemcpyKind kind);) je patrné, že oproti standardní funkci memcpy je tu ješte enumerátor cudaMemcpyKind. Ten může nabývat hodnot cudaMemcpyHos­tToHost, cudaMemcpyHos­tToDevice, cudaMemcpyDevi­ceToHost,  cudaMemcpyDevi­ceToDevice.

cudaMemcpyKind je popsán v tabulce č. 3.2.

Enumerátor cudaMemcpyKind popis
cudaMemcpyHos­tToHost Kopírování paměti z hosta do paměti hosta
cudaMemcpyHos­tToDevice Kopírování paměti z hosta do globální paměti
cudaMemcpyDevi­ceToHost Kopírování paměti z globální paměti do paměti hosta
cudaMemcpyDevi­ceToDevice Kopírování paměti z globální paměti do paměti zařízení

Tab. 3.3 Bližší popis cudaMemcpyKind

V programu 3.1. numbers.cu konkrétně v kernelu numbers je výpočet indexu threadu odlišný než u hello_cuda.cu a to takto:

int idx = blockDim.x * blockIdx.x + threadIdx.x;

Je tomu tak proto, že již nepracujeme pouze s jediným blokem o 12 vláknech, ale s celou řadou bloků. Ve funkci main jsme definovali, že velikost pole bude num_size = 128512, kde počet vláken na blok je definován proměnnou.

size_t num_threads_per_block = 128;


Z toho jsme pak spočítali požadovanou velikost mřížky (celková velikost pole děleno počtem vláken na blok), tedy:

num_blocks = num_size/num_threads_per_block;

Z výše popsaného rovněž vyplívá, že každých 128 vláken může být synchronizováno (v kernelu pomocí __syncthreads()) a mohou přistupovat ke společné sdílené paměti. A jak jsme dospěli ke vzorci blockDim.x * blockIdx.x + threadIdx.x a co znamená blockDim a blockIdx? Obrázek 3.1 nám bude k tomu nápomocen (pro lepší pochopení uspořádání dvoudimenzionální mřížky, bloků a
vláken můžeme využít i obrázek 2.1 z předcho­zího dílu).

Cuda 3.1

Obr. 3.1 Výpočet idx

Kromě built-in proměnné threadIdx existují ještě proměnné blockDim a blockIdx. Jak vidíme z obrázku 3.1 ke vzorci idx = blockDim.xbloc­kIdx.x + threadIdx.x jsme dospěli tak, že jsme vynásobili počet bloků (blockDim) s aktuálním blockem (blockIdx) a k němu přičetli index vlákna.

Z toho plyne, že threadIdx je absolutní k bloku, nikoliv k celé mřížce (celkovému počtu bloků). Idx nám pak slouží k přístupu do globální paměti, jak vstupního, tak výstupního pole čísel:

out[idx] = in[idx]/2;

Jinými slovy, každé vlákno zpracovává jedno číslo z vstupního pole in o indexu idx a ukládáho do výstupního pole out o indexu idx. Zaměřme se nyná na obrázek 3.2, kde jsou vidět dva příklady stejné mřížky. Ta je jinak “mapovaná” na MP na zařízení. V horní části obrázku vidíme mřížku o velikosti (2, 4) “mapovanou” na 2 MP. V
dolní části obrázku vidíme stejnou mřížku (2, 4), která je mapována na 4 MP. 

Znamená to tedy, že nemusíme nijak modifikovat náš zdrojový kód pro různá hw uspořádání bloků v MP.

Cuda 3.2


Obr. 3.2 Dva příklady uspořádání bloků v MP

Sdílená paměť


Sdílená paměť se alokuje přímo v execution configuration, nap­ř:

<<<pocet_bloku, velikost_bloku, velikost_sdilene_pameti_v_bytech>>>

Jak je vidět, velikost sdílené paměti určuje 3. parametr v <<<>>> a dále ji musíme definovat v kernelu, například takto:

__global__ void my_kernel () {
       //definice s_data jako sdilena pamet (velikost je definovana v execution configuration)
       extern __shared__ int s_data[];
}


Sdílená paměť definovaná v kernelu musí mít identifikátor __shared__ a poté pokračujeme typem proměnné (int, float, apod.) jak jsme zvyklí z jazyka C. Klíčové slovo extern je uvedeno proto, že velikost této paměti neurčujeme v kernelu, ale v execution configuration. Existuje možnost definovat sdílenou paměť přímo v kernelu, a to následovně:

__global__ void my_kernel () {
       //definice s_data jako sdilena pamet
       __shared__ int s_data[velikost_v_bytech];
}


Chtěl bych upozornit na fakt, že přístup do sdílené paměti tak, jak nám ukazuje obr. 1.3 (první díl) a ještě připomíná tabulka 3.1, je v rámci bloku. To tedy znamená, že pokud budeme alokovat sdílenou paměť, musíme její velikost počítat v rámci jednoho bloku, nikoliv celé mřížky. Například když se podíváme na program 3.1, vidíme že pro pole čísel jsme alokovali 128* 512 integerů a rovněž jsme
definovali, že počet vláken na blok bude 128. Pokud v tomto případě budeme pracovat se sdílenou pamětí, její velikost nebude 128* 512 integerů, ale jen 128 integerů. Je tomu tak proto, že v jednom bloku je 128 vláken (připomínám, že jedno číslo zpracovává jedno vlákno ⇒ jeden integer na jedno
vlákno * počet vláken v bloku). Upravený zdrojový kód programu se sdílenou pamětí:

#include <stdio.h>
__global__ void numbers (int* out, int* in) {
       //definujeme promennou s_data jako sdilenou pamet
       extern __shared__ int s_data[];
       // spocitame index do pole in a out pro konkretni vlakno
       // pripominam, ze pracujeme pouze s jednorozmernym polem,
       // tudiz jsme definovali x-slozku velikosti mrizky a bloku v execution configuration.
       int idx = blockDim.x * blockIdx.x + threadIdx.x;
       //delime cislo in[idx] dvema a ukladame jej do sdilene pameti
       s_data[threadIdx.x] = in[idx]/2;
       //blokuj dokud tohoto bodu nedosahnou vsechny vlakna
       __syncthreads ();
       // ukladame do vystupniho pole
       // ma smysl, pokud idx nepresahne velikost vstupniho (a vystupniho) pole
       out[idx] = s_data[threadIdx.x];
}
int main (int argc, char** argv) {
       //pomocna promenna
       bool result = true;
       //ukazatel na pole integeru pro praci s cisly v kernelu
       int* num_h;
       //ukazatel na pole int v globalni pameti
       int* num_d;
//ukazatel na vystupni pole v globalni pameti
int* num_out_d;
//pozadovana velikost pole
size_t num_size = 128*512;
//pocet vlaken na jeden blok
int num_threads_per_block = 128;
//velikost mrizky
int num_blocks = num_size/num_threads_per_block;
//pozadovana velikost pole v bytech
size_t num_size_bytes = sizeof (int)*num_size;
//alokace pameti pole
num_h = (int*)malloc (num_size_bytes);
//alokujeme pole num_d v globalni pameti
cudaMalloc ((void**) &num_d, num_size_bytes);
//alokujeme vystupni pole num_out_d v globalni pameti
cudaMalloc ((void**) &num_out_d, num_size_bytes);
//naplnime pole cisly
for (unsigned int i = 0; i < num_size; i++) {
        num_h[i] = i*2;
}
//kopirovani pole int z hosta do globalni pameti
cudaMemcpy (num_d, num_h, num_size_bytes, cudaMemcpyHostToDevice);
//spocitame velikost sdilene pameti (pocet vlaken v jednom bloku * velikost integeru)
size_t shared_mem_size = num_threads_per_block*sizeof(int);
//volame kernel
numbers<<<num_blocks, num_threads_per_block, shared_mem_size>>> (num_out_d,
num_d);
       //cekej, dokud vsechny vlakna neskonci
       cudaThreadSynchronize();
       //nyni muzeme zkopirovat vystupni data num_out_d do num_h
       cudaMemcpy (num_h, num_out_d, num_size_bytes, cudaMemcpyDeviceToHost);
       //zkontrolujeme, zda vysledek je spravny
       for (unsigned int i = 0; i < num_size; i++) {
               if (i != num_h[i]) {
                        result = false;
                        printf ("Vysledek je spatny! %d\n", i);
                        break;
               }
       }
       //pokud je vysledek spravny, vytiskneme oznameni
       if (result) {
               printf ("Vysledek je spravny!\n");
       }
       //uvolneni prostredku globalni pameti
       cudaFree (num_d);
       cudaFree (num_out_d);
       //uvolneni prostredku v pameti hosta
       free (num_h);
       return 0;
}

Kód 3.2 Program numbers_shared.cu se sdíleným přístupem do paměti

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

Jelikož práce s globální pamětí je neefektivní, protože je nekešovaná (a také proto, že není na chipu, ale v DRAM zařízení) a tudíž pomalá, existuje technika – sdílený přístup do paměti (coalesced memory access), pomocí které můžeme výrazně urychlit přístup do této paměti.

Každá instrukce na MP se zpracovává po skupinách nazvaných warp. Warp je velikosti 32, tedy 32 vláken se zpracovává současně. Half-warp je první nebo druhá půlka warpu, tedy buď 0–15 vlákno nebo 16–31. vlákno. Jedná se o transakci, kde 16 vláken (half-warp) přistupuje současně do globální paměti (tak se přístup do ní “jakoby zrychlí”).


Pro sdílený přístup do paměti je nutné, aby přístup do globální paměti splňoval následující podmínky:

1. Velikost paměťového elementu, ke které přistupujeme, je 4, 8 nebo 16 bytů (např. int, float, ale už ne char!).
2. Vlákna k elementům přistupují sekvenčně: tedy k n-tému elementu pouze n-té vlákno (nepřistupují-li všechna vlákna do paměti, jedná se o tzv. divergent warp).
3. Všech 16 elementů leží ve stejném segmentu, přičemž adresa prvního elementu musí být zarovnána k 16násobku velikosti elementu.

Pokud jsou výše uvedené podmínky splněny, jedná se o sdílený přístup do paměti (16 vláken souběžně přistupuje do globální paměti). V opačném případě je transakce zúžena pouze na jediné vlákno (v jeden okamžik přistupuje pouze jedno vlákno do sdílené paměti)! Rozdíl uvedených transakcí je samozřejmě markantní a proto se vyplatí dodržovat výše popsaná pravidla. Pro lepší názornost sdíleného přístupu nám poslouží obrázek 3.3.

Cuda 3.3

Obr. 3.3 Příklad sdíleného přístupu do paměti

Cuda 3.4

Obr. 3.4 Příklad nesdíleného přístupu do paměti

Závěr


V dnešním díle jsme se naučili, jak pracovat se sdílenou pamětí a jak optimalizovat přístup do globální paměti. Příště dokončíme sdružený přístup do paměti a zaměříme se na slibované debuggování, které se už do dnešního dílu nevešlo.

Použitá literatura:
[1] CUDA 2.2 Programming Guide
[2] CUDA 2.2 QuickStar­t Guide

UX DAy - tip 2

Použité obrázky byly převzaty z těchto odkazů:

[1] https://visualization.hpc.mil/…amming_Model
[2] NVIDIA_CUDA_Tu­torial_No_NDA_A­pr08.pdf

Byl pro vás článek přínosný?