CUDA: více o sdruženém přístupu do globální paměti

Lukáš Zaorálek 10. 8. 2009

Předchozí třetí díl seriálu o technologii CUDA nás naučil, jak pracovat se sdílenou pamětí a jak optimalizovat přístup do pomalé globální paměti. Využili jsme při tom sdruženého přístupu do paměti, o kterém si v dnešním díle povíme více, a dále probereme debuggování CUDA aplikací.

Rekapitulace

Již víme, jak optimalizovat přístup do globální paměti pomocí tzv. sdruženého přístupu do globální paměti (coalesced memory access). Sdružený přistup do paměti znamená, že 16 vláken (tzv. half-warp) současně přistupuje do globální paměti a tím se zakryje latence přístupu do paměti. V opačném případě (nesdružený přístup do paměti) by přistupovalo do globální paměti jen jedno vlákno v jeden okamžik. Aby se jednalo o sdružený přístup do globální paměti, je nutné splnit několik podmínek:

    1. Velikost paměťového elementu, ke kterému 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.

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


Z výše uvedených podmínek je patrné, že akceptovatelná velikost proměnné jsou 4, 8 nebo 16 bytů. Proto bude vyhovovat pro sdružený přístup do paměti typ proměnné float nebo int (oba typy jsou zarovnané na 4 byty). Pokud se dále podíváme na vestavěné typy CUDA jako je int{2–4} nebo float{2–4} zjistíme, že pro 1. podmínku uvedenou výše nám nevyhovuje float3 (3×4 se nerovná 4, 8 nebo 16) nebo int3 a rovněž obyčejný char, který má 1 byte. Jak se zachovat v těchto případech?


Ukažme si na jednoduchém příkladu. Mějme pole typu float3 o N prvcích, kde ke každému prvku přičteme hodnotu value. Zdrojový kód takového programu může vypadat takto:

#include <stdio.h>
#include <assert.h>
__global__ void noncoalesced_float3_kernel (float3* out, float3* in, float value) {
       //index vlakna v ramci vsech bloku
       int idx = blockIdx.x * blockDim.x + threadIdx.x;
       //prace s cislem typu float3
       in[idx].x += value;
       in[idx].y += value;
       in[idx].z += value;
       //ulozime element do globalni pameti do vystupniho pole
       out[idx] = in[idx];
}
int main (int argc, char** argv) {
       //pole pro praci s float3
       float3 *array_in_h, *array_in_d, *array_out_h, *array_out_d;
       //pomocna promenna
       float value;
       //velikost pole (pocet float3 elementu v poli)
       size_t array_size = 8*64;
       //velikost pole v bytech
       size_t array_size_bytes = array_size * sizeof(float3);
       //pocet vlaken na jeden blok
       int num_threads_per_block = 64;
       //pocet bloku
       int num_blocks = array_size/num_threads_per_block;
       //alokace pozadovane pameti pro pole typu float3
       array_in_h = (float3*)malloc (array_size_bytes);
       //naplneni pole cisly
       for (unsigned int idx = 0; idx < array_size; idx++) {
                array_in_h[idx].x = (float)idx;
                array_in_h[idx].y = (float)idx;
                array_in_h[idx].z = (float)idx;
       }
       //alokace vystupniho a vstupniho pole v globalni pameti
       cudaMalloc ((void**) &array_out_d, array_size_bytes);
       cudaMalloc ((void**) &array_in_d, array_size_bytes);
       //kopirovani naplneneho pole z hosta do globalni pameti zarizeni
       cudaMemcpy (array_in_d, array_in_h, array_size_bytes, cudaMemcpyHostToDevice);
       //hodnota, ktera se bude pricitat k jednotlivym slozkam (x, y, z) elementu v poli typu float3
       value = 3.0f;
       //volani kernelu
       noncoalesced_float3_kernel<<<num_blocks, num_threads_per_block="">>> (array_out_d,
array_in_d, value);
       //pockej dokud nezkonci vsechny vlakna
       cudaThreadSynchronize ();
       //alokace pameti pro vystupni pole (v hostu)
       array_out_h = (float3*)malloc (array_size_bytes);
       //kopirovani vystupniho pole v globalni pameti do vystupniho pole v pameti hosta
       cudaMemcpy (array_out_h, array_out_d, array_size_bytes, cudaMemcpyDeviceToHost);
       //kontrola, zda soucet je v poradku
       for (unsigned int idx = 0; idx < array_size; idx++) {
               assert (array_out_h[idx].x == idx+value);
               assert (array_out_h[idx].y == idx+value);
               assert (array_out_h[idx].z == idx+value);
       }
       //pokud ano, vypise se zprava nize:
       printf ("Ok!\n");
       //uvolneni alokovane pameti vstupniho a vystupniho pole v pameti hosta
       free (array_in_h);
  free (array_out_h);
  //uvolneni alokovane pameti vstupniho a vystupniho pole v globalni pameti zarizeni
  cudaFree (array_in_d);
  cudaFree (array_out_d);
  //ukonceni programu
  return 0;
}

Kód 4.1: Soubor nocoalesced.cu nesdružený přístup do paměti pole typu float3

Když si graficky znázorníme přístup do paměti, obr. 4.1, uvidíme, že se skutečně nejedná o sdružený přístup do paměti. Každé vlákno (označené jako t{1,2,3}) přistupuje do paměti složek float3 x, y a z. Jinými slovy, každé vlákno nepřistupuje kontinuálně ke složkám x, y a z. Řešením je buď použití float4, kde nebudeme používat čtvrtou složku a nebo použití sdílené paměti, o které byla řeč ve 3. díle. Ukážeme si druhý ”těžší” způsob. Musíme náš kód 4.1 trochu pozměnit.

Obr 4.1

Obr. 4.1: Nesdružený přístup do paměti pole typu float3

Scénář změn kódu bude rozdělen do několika jednoduchých kroků:

1. Abychom mohli využít sdružený přístup do paměti, je nutné pracovat s polem float, nikoliv s float3. Předpis našeho kernelu bude vypadat takto:

__global__ void coalesced_float3_kernel (float* out, float* in, float value);
2. Ve funkci main provedeme jen přetypování z pole float3* na pole float* při volání kernelu (zbytek kódu se ve funkci main nemění):
    coalesced_float3_kernel<<<num_blocks,num_threads_per_block,shared_mem_size>>>
        ((float*)array_out_d, (float*)array_in_d, value);

3. Klíčovým krokem je sdružený přístup do globální paměti a uložení hodnot do paměti sdílené:
   s_data[threadIdx.x] = in[index];
   s_data[threadIdx.x+64] = in[index+64];
   s_data[threadIdx.x+128] = in[index+128];
   Je dodržený kontinuální přístup, kde se přistupuje buď ke složce x, nebo ke složce y nebo ke složce z, ale ne k x,y,z současně, tak jako v kódu 4.1.
4. Při práci se samostatným číslem float3 je nutné toto číslo “dostat” z pole float:
   float3 number = ((float3)s_data)[threadIdx.x];
   A poté ho zpět uložit do sdílené paměti:
   ((float3*)s_data)[threadIdx.x] = number;
5. Nakonec uložíme čísla do výstupního pole v globální paměti. Přístup do paměti je obdobný jako v kroku 3:
   out[index] = s_data[threadIdx.x];
   out[index+64] = s_data[threadIdx.x+64];
   out[index+128] = s_data[threadIdx.x+128];

Nesmíme ovšem zapomenout na synchronizaci vláken (v rámci bloku) před a po kroku 4. Upravený kód 4.2 využívající sdruženého přístupu do globální paměti pole typu float3:

#include <stdio.h>
#include <assert.h>
__global__ void coalesced_float3_kernel (float* out, float* in, float value) {
       //sdilena pamět, velikost je definovana ve funkci main
       extern __shared__ float s_data[];
       //index vlakna v ramci vsech bloku
       int index = blockIdx.x * blockDim.x + threadIdx.x;
       //ulozime jednotlive slozky float3 do sdilene pameti
       //v s_data je pouzit threadIdx, protoze se jedna o pamet v ramci jednoho bloku
       //sdruzeny pristup do globalni pameti
       s_data[threadIdx.x] = in[index];
       s_data[threadIdx.x+64] = in[index+64];
       s_data[threadIdx.x+128] = in[index+128];
       //nez budeme pokracovat dale, je nutne vsechny zkopirovat z globalni pameti do sdilene
       __syncthreads();
       //pretypujeme pozadovane cislo na float3
       float3 number = ((float3*)s_data)[threadIdx.x];
       //prace s cislem typu float3
       number.x += value;
       number.y += value;
       number.z += value;
       //pretypovani cisla typu float3 zpet na float* pole
       ((float3*)s_data)[threadIdx.x] = number;
       //pockame dokud neni prace s cislem dokoncena u vsech vlaken
       __syncthreads();
       //vse prekopirujeme ze sdilene pameti na vystup do globalni pameti
       //sdruzeny pristup do globalni pameti
       out[index] = s_data[threadIdx.x];
       out[index+64] = s_data[threadIdx.x+64];
       out[index+128] = s_data[threadIdx.x+128];
}
int main (int argc, char** argv) {
       //pole pro praci s float3
       float3 *array_in_h, *array_in_d, *array_out_h, *array_out_d;
       //pomocna promenna
       float value;
       //velikost pole (pocet float3 elementu v poli)
       size_t array_size = 8*64;
       //velikost pole v bytech
       size_t array_size_bytes = array_size * sizeof(float3);
       //pocet vlaken na jeden blok
       int num_threads_per_block = 64;
       //pocet bloku
       int num_blocks = array_size/num_threads_per_block;
       //velikost sdilene pameti v bytech (sdilena pamet je pouze v ramci jednoho bloku!)
       int shared_mem_size = num_threads_per_block * sizeof (float3);
        //alokace pozadovane pameti pro pole typu float3
        array_in_h = (float3*)malloc (array_size_bytes);
        //naplneni pole cisly
        for (unsigned int idx = 0; idx < array_size; idx++) {
                array_in_h[idx].x = (float)idx;
                array_in_h[idx].y = (float)idx;
                array_in_h[idx].z = (float)idx;
        }
        //alokace vystupniho a vstupniho pole v globalni pameti
        cudaMalloc ((void**) &array_out_d, array_size_bytes);
        cudaMalloc ((void**) &array_in_d, array_size_bytes);
        //kopirovani naplneneho pole z hosta do globalni pameti zarizeni
        cudaMemcpy (array_in_d, array_in_h, array_size_bytes, cudaMemcpyHostToDevice);
        //hodnota, ktera se bude pricitat k jednotlivym slozkam (x, y, z) elementu v poli typu float3
        value = 3.0f;
        //volani kernelu
        coalesced_float3_kernel<<<num_blocks, num_threads_per_block, shared_mem_size>>>
((float*)array_out_d, (float*)array_in_d, value);
        //pockame dokud neskonci vsechny vlakna
        cudaThreadSynchronize ();
        //alokace pameti pro vystupni pole (v hostu)
        array_out_h = (float3*)malloc (array_size_bytes);
        //kopirovani vystupniho pole v globalni pameti do vystupniho pole v pameti hosta
        cudaMemcpy (array_out_h, array_out_d, array_size_bytes, cudaMemcpyDeviceToHost);
        //kontrola, zda soucet je v poradku
        for (unsigned int idx = 0; idx < array_size; idx++) {
                assert (array_out_h[idx].x == idx+value);
                assert (array_out_h[idx].y == idx+value);
        assert (array_out_h[idx].z == idx+value);
}
//pokud ano, vypise se zprava nize:
printf ("Ok!\n");
//uvolneni alokovane pameti vstupniho a vystupniho pole v pameti hosta
free (array_in_h);
free (array_out_h);
//uvolneni alokovane pameti vstupniho a vystupniho pole v globalni pameti zarizeni
cudaFree (array_in_d);
cudaFree (array_out_d);
//ukonceni programu
return 0;
}

Kód 4.2: Soubor coalesced.cu sdružený přístup do globální paměti pole typu float3

Možná přemýšlíte, jak využít sdruženého přístupu do globální paměti, a to například pro pole 1bytových elementů jako je pole typu char* (kde jasně porušujeme 1. podmínku uvedenou výše). Kompletní zdrojový kód řešení nechám jako domácí úkol, jen napovím, že je proto možné využít typ (unsigned) int. Typ int má velikost 4bytů, a proto se do něj “vejdou” 4 chary. Pomocné metody pro uložení či nahrání charů z, resp. do jednoho int může vypadat takto:

unsigned int store_chars (char* str) {
       return (str[3] << 24)|(str[2] << 16)|(str[1] << 8)|(str[0] << 0);
}
char* load_chars (unsigned int num) {
       char* str = (char*)malloc (sizeof (char)*4);
       str[3] = num >> 24; str[2] = num >> 16; str[1] = num >> 8; str[0] = num >> 0;
       return str;
}

Kód 4.3: Pomocné funkce pro uložení/nahrání 4 charů z/do int

Pro struktury, které nejsou zarovnané na velikost 4, 8 nebo 16 bytů použíjeme jednu z následujících možností:

1. výše uvedeného způsobu sdruženého přístupu do paměti;
2. strukturu polí (structure of arrays, zkratka SoA). Nedoporučuje se pole struktur (array of struktures, zkratka AoS), protože může dojít ke stejným problémům jako u pole typu float3 (rovněž struktura) zobrazující obr. 4.1.
3. zarovnání struktur užitím __align__ (x), kde x bude 4, 8 nebo 16 bytů, např:

struct __align(8)__ {
    float a;
    float b;
    };

nebo

struct __align(16)__ {
float a;
float b;
float c;
float d;
};

Shrnutí přístupu do globální paměti

Tam, kde je to možné, raději použijeme sdílenou paměť než globální paměť. Globální paměť má latenci 400 – 600 cyklů a je velmi neefektivní pro přístup do ní. Pokud už musíme použít globální paměť (například pro vstupní nebo výstupní data), snažíme se zakrýt latenci využitím sdruženého přístupu do této globální paměti (coalesced memory access).

Debuggování CUDA programů

Debuggování programů se provádí pomocí programu cuda-gdb, který je podobný známému debuggeru gdb. Abychom mohli debuggovat CUDA program, musíme nejdříve zkompilovat zdrojové kódy kompilerem nvcc s parametry -g -G. První parametr -g znamená debug pro kód hosta. Druhý parametr -G znamená debug pro kód zařízení (tedy kernelů). Například:

$ nvcc noncoalesced.cu -g -G -o noncoalesced.

Spuštění cuda-gdb:

$ cuda-gdb noncoalesced

Až budeme spouštět cuda-gdb, musíme si dát pozor na to, abychom měli vypnutá Xka. Je tomu tak kvůli ochraně. Špatně napsaný CUDA program by tak mohl způsobit zatuhnutí nebo pád Xek. Pokud budeme chtít debuggovat aplikaci, používáme standardní ovládání jako u gdb.

widgety

Závěr

Dnešní článek dokončil povídání o sdruženém přístupu do paměti. Ukázali jsme si základy debuggování CUDA programů. V příštím díle si ukážeme praktickou ukázku debuggování a povíme si o bankách sdílené paměti.

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?
Lupa.cz: Kde leží hardwarový pupek světa?

Kde leží hardwarový pupek světa?

Podnikatel.cz: „Lex Babiš“ Babišovi paradoxně pomůže

„Lex Babiš“ Babišovi paradoxně pomůže

Lupa.cz: Blíží se konec Wi-Fi sítí bez hesla?

Blíží se konec Wi-Fi sítí bez hesla?

Vitalia.cz: Jak Ondra o astma přišel

Jak Ondra o astma přišel

Podnikatel.cz: ČSSZ posílá přehled o důchodovém kontě

ČSSZ posílá přehled o důchodovém kontě

Vitalia.cz: 5 chyb, které děláme při skladování potravin

5 chyb, které děláme při skladování potravin

Podnikatel.cz: Nemá dluhy? Zjistíte to na poště

Nemá dluhy? Zjistíte to na poště

Lupa.cz: Další Češi si nechali vložit do těla čip

Další Češi si nechali vložit do těla čip

Podnikatel.cz: Tyto pojmy k #EET byste měli znát

Tyto pojmy k #EET byste měli znát

Podnikatel.cz: Byla finanční manažerka, teď cvičí jógu

Byla finanční manažerka, teď cvičí jógu

DigiZone.cz: Wimbledon na Nova Sport až do 2019

Wimbledon na Nova Sport až do 2019

DigiZone.cz: Digi Slovakia zařazuje stanice SPI

Digi Slovakia zařazuje stanice SPI

DigiZone.cz: Mordparta: trochu podchlazený 87. revír

Mordparta: trochu podchlazený 87. revír

DigiZone.cz: Světový pohár v přímém přenosu na ČT

Světový pohár v přímém přenosu na ČT

DigiZone.cz: Parlamentní listy: kde končí PR...

Parlamentní listy: kde končí PR...

Lupa.cz: Jak se prodává firma za miliardu?

Jak se prodává firma za miliardu?

Lupa.cz: Patička e-mailu závazná jako vlastnoruční podpis?

Patička e-mailu závazná jako vlastnoruční podpis?

Vitalia.cz: Antibakteriální mýdla nepomáhají, spíš škodí

Antibakteriální mýdla nepomáhají, spíš škodí

Root.cz: Hořící telefon Samsung Note 7 zapálil auto

Hořící telefon Samsung Note 7 zapálil auto

Vitalia.cz: Voda z Vltavy před a po úpravě na pitnou

Voda z Vltavy před a po úpravě na pitnou