Jak programovat CUDA ooooooooooooooooooooooooo
CUDA, nástroje a knihovny
Jiří Matela
podzim 2009
□           SP             -             =            -^       ^o^O
EBB
ni
Jak programovat CUDA ooooooooooooooooooooooooo
Proč programovat GPU
□          SP           -           =           -^      ^o^o
■I
ni
Rekapitulace
What if you could go from S F to NYC ...in 3 minutes?	-•I           »Ä	
	W-  P	
	IG            ^	
		
What does it mean to someone who cares how long it takes to do something when you
can speed things up 140 times, 100 times or even 50 times? It is like being able to go
from San Francisco to New York in three minutes. A speed up of that kind is
transformative. It would completely transform adjacent industries.
- Jen-Hsun Huang, nVidia CEO < í ►
Jiří Matela           CUDA, nástroje a knihovny
Rekapitulace
o»o
Jak programovat CUDA                             NVCC
ooooooooooooooooooooooooo        ooooooo
Knihovny
ooo
Rekapitulace
What does it mean to someone who cares how long it takes to do something when you
can speed things up 140 times, 100 times or even 50 times? It is like being able to go
from San Francisco to New York in three minutes. A speed up of that kind is
transformative. It would completely transform adjacent industries.
— Jen-Hsun Huang, nVidia CEO
^)<\(j*.
Jak programovat CUDA ooooooooooooooooooooooooo
• Proč programovat GPU
s
■O Q-C^
EBB
Jak programovat CUDA ooooooooooooooooooooooooo
•  Proč programovat GPU
•  GPU architektura (vs. CPU)
s
■O Q-C^
EBB
Jak programovat CUDA
ooooooooooooooooooooooooo
Knihovny
ooo
Rekapitulace
•  Proč programovat GPU
•  GPU architektura (vs. CPU)
•  CUDA (Compute Unified Device Architecture)
Architektura CUDA
CUDA Architecture
□      s       ~       =
CUDA, nástroje a knihovny
Rekapitulace oo«
Jak programovat CUDA ooooooooooooooooooooooo
Rekapitulace
•  Proč programovat GPU
•  GPU architektura (vs. CPU)
9 CUDA (Compute Unified Device Architecture)
Hierarchie vláken
Grid Block(0, 0)	Block (1,0)	Block (2,0)
Block (O,!)''	Block (1,1)	-Block (2,1)
Thread (3, 2)
□ ►   < ť5> ►   < -= ►   < = ►
CUDA, nástroje a knihovny
Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOi
Rekapitulace
•   Proč programovat GPU
•   GPU architektura (vs. CPU)
•   CUDA (Compute Unified Device Architecture)
Hierarchie pamětí
Block (O, O)      Block (1, O)      Block Í2, O)
Block (O, 1)      Block (1, 1)      Block (2, 1)
;(0,0)               Block U. 0)
Block í O, I!               Bloch U. 1)
&
■0 0,0
CUDA, nástroje a knihovny
Jak programovat CUDA
Proč programovat GPU GPU architektura (vs. CPU)
CUDA (Compute Unified Device Architecture)
• Dvě API
■I
S
•f) <\(y
Jak programovat CUDA
•  Proč programovat GPU
•  GPU architektura (vs. CPU)
a CUDA (Compute Unified Device Architecture)
•  Dvě API
•  Ukázkový příklad
•   Syntaktická rozšíření jazyka C - např.: __device__
•   Volání runtime API -např.: cudaMallocO
■I
S
•f) <\(y
Jak programovat CUDA
•OOOOOOOOOOOOOOOOOOOOOOCH
Runtime API vs. Driver API
Vytvářet CUDA aplikace lze užitím buďto Runtime API nebo Driver API.
CPU
Aplikace
CUDA Runtime API
Aplikace
CUDA Driver API
□ ► < s
Jiří Matela           CUDA, nástroje a knihovny
►       -š      -00,0
Jak programovat CUDA
OOOO   DOOOOOOOOOOOOOOOO
Možnosti rozhraní
Rozhraní umožňují provádět na úrovni hostitelského systému (kód vykonávaný na CPU) následující operace
•  Správa zařízení
•  Práce s kontextem
•   Práce s moduly
•   Konfigurace výpočtu
•  Paměťové operace
•  Práce s texturami
•  Spolupráce s OpenGL a Direct3D
EBB
S
•f)<\(y
•   Runtime API a C for ČUDA - množina rozšírení jazyka C
•  Automatická inicializace, práce s kontextem a práce s kernely
•   Konfigurace výpočtu (volání kernelu) - syntaktický konstrukt (rozšíření jazyka C)
•   Kód používající rozšíření musí být přeložen nvcc kompilátorem
•  Jinak lze hostitelský kód přeložit pomoci gcc
Rekapitulace
Jak programovat CUDA
Příklad kódu používajícího CUDA rozšíření jazyka C
Konfigurace CUDA kernelu addvecO
int   mainQ    {
addvec<«N/BLOCK ,    BL0CK»>(d_a,    d_b ,    d_c
}
■I
S
•f)<\(y
Rekapitulace
Jak programovat CUDA
Příklad kódu používajícího CUDA rozšíření jazyka C
Konfigurace CUDA kernelu addvecO
int   mainQ    {
addvec<«N/BLOCK ,    BL0CK»>(d_a,    d_b ,    d_c
}
Překlad:
$ nvcc -I/usr/local/cuda/include -L/usr/local/cuda/lib \ -lcudart -o vecadd vecadd.cu
□ ► 4 s ► <■=► < ■= ►
■O q,C^
■I
Knihovna runtime API:
•  cuda_runtime_api.h - nízkoúrovňové C funkce
•  cuda_runtime.h - vysokoúrovňové C++ funkce - obaluje C api
•  Funkce pro alokaci a dealokaci paměti
•  Přenos dat z/do karty
•  Správa karet - výběr a konfigurace karty
•  Podpora emulace karty na CPU - debuging
•  Volání prefixováné cuda*
□          S           ~           =          -^"O^O
Rekapitulace
Jak programovat CUDA
Přiklad kódu používajícího runtime API volaní
Informace o kartě
int   mainQ    {
cudaGetDeviceCount (&devCount);
printf("Available    devices:   %d\n",   devCount
cudaGetDeviceProperties(devProp,    0);
printf (" Device :   %d\n" ,    i );
printf(" Name:   %s\n" ,    devProp—>name);
■I
S
•f)<\(y
Rekapitulace
Jak programovat CUDA
Přiklad kódu používajícího runtime API volaní
Informace o kartě
int   mainQ    {
cudaGetDeviceCount (&devCount);
printf("Available    devices:   %d\n",   devCount
cudaGetDevieeProperties(devProp, 0);
printf (" Device : %d\n" , i );
printf(" Name: %s\n" , devProp—>name);
}
Překlad:
$ gcc -I/usr/local/cuda/include -L/usr/local/cuda/lib \ -leudart -x c -o info info.cu
□ ► 4 s ► <■=► < ■= ►
■O q,C^
■I
Nízkoúrovňové rozhraní pro programování CUDA aplikací. (V pomyslné hierarchii je položeno níž než runtime API)
•  Více kontroly nad kartami -jedno CPU vlákno může pracovat s více kartami
•  Neobsahuje žádné rozšíření jazyka C
•  Umožňuje pracovat s binárním kódem a assemblerem (PTX)
•  Složitější programování, upovídanější syntax
•  Složitější debugging
•  Žádná GPU emulace
•  Volání prefixováné cu*
Prostředí CUDA výpočtu představuje context (analogie CPU procesu nebo vlákna).
•  Vztahuje se ke konkrétnímu GPU zařízení
•  Zastřešuje všechny zdroje a vykonané akce
•  Má vlastní 32-bit paměťový prostor (paměťové ukazatele nelze mezi kontexty přenášet)
•  CPU vlákno může v danou chvíli používat vždy jen jeden kontext
•  Kontexty lze mezi vlákny předávat
•  Použití více karet jedním CPU vláknem
Jak programovat CUDA
Příklad inicializace contextu
Inicializace kontextu je v případě runtime API implicitní, zatímco v případě driver API vyžaduje několik příkazů:
CUcontext   cont; CUdevice   dev;
culnit(0);    //   0   je    povinná ,    parametr   zatím   nemá   význam cuDeviceGet (&cont ,    0);    //   vyber    první    kartu    (0) cuCtxCreate(&cont,    CU_CTX_SCHED_AUTO,    dev));
■I
S
•f) <\(y
Jak programovat CUDA
S kernely se pracuje jako s moduly, které jsou (obdobně jako GLSL shadery) nahrávány za běhu.
•  Binární moduly - kompilovány pro konkrétní architekturu, mohou být pomalejší nebo nekompatibilní na budoucích architekturách
•  PTX moduly - kompilovány až v době natažení (PTX je meta assembler, jehož instrukce jsou nejprve přeloženy do skutečné instrukční sady dané architektury a následně pak do binárního kódu)
Jak programovat CUDA
Příklad konfigurace a spuštění modulu
CUmodule   myModule; CUfunction   myKern;
//nahrát    modul,    získat    kernel    "addvec" cuModuleLoad(&myModule,    "vectorAdd.cubin"); cuModuleGetFunction(&myKern,    myModule,    "addvec"
EBB
S
•f) <\(y
Jak programovat CUDA
Příklad konfigurace a spuštění modulu
CUmodule   myModule; CUfunction   myKern;
//nahrát    modul,    získat    kernel    "addvec" cuModuleLoad(&myModule,    "vectorAdd.cubin"); cuModuleGetFunction(&myKern,    myModule,    "addvec"
//    nastavit    parametry   thread    bloku cuFuncSetBlockShape(myKern,    x,    y,    z);
■I
S
•f) <\(y
Jak programovat CUDA
Příklad konfigurace a spuštění modulu
CUmodule   myModule; CUfunction   myKern;
//nahrát    modul,    získat    kernel    "addvec" cuModuleLoad(&myModule,    "vectorAdd.cubin"); cuModuleGetFunction(&myKern,    myModule,    "addvec");
//    nastavit    parametry   thread    bloku cuFuncSetBlockShape(myKern,    x,    y,    z);
//    nakopírovat    size    bajtů    z   &ptr   do    prostoru    parametrů //    kernelu    myKern    na    pozici    offset cuParamSetv(myKern,    offset,   &ptr,    size);
EBB
S
•f)<\(y
Jak programovat CUDA
Příklad konfigurace a spuštění modulu
CUmodule   myModule; CUfunction   myKern;
//nahrát    modul,    získat    kernel   "addvec" cuModuleLoad(&myModule,    "vectorAdd.cubin"); cuModuleGetFunction(&myKern,    myModule,    "addvec");
//    nastavit    parametry   thread    bloku cuFuncSetBlockShape(myKern,    x,    y,    z);
//    nakopírovat    size    bajtů    z  &ptr   do    prostoru    parametrů //    kernelu    myKern   na    pozici    offset cuParamSetv(myKern,    offset,   &ptr,    size);
//   celková    velikost    parametrů cuParamSetSize(myKern,    offset);
■I
S
•f)<\(y
Jak programovat CUDA
Příklad konfigurace a spuštění modulu
CUmodule   myModule; CUfunction   myKern;
//nahrát    modul,    získat    kernel   "addvec" cuModuleLoad(&myModule,    "vectorAdd.cubin"); cuModuleGetFunction(&myKern,    myModule,    "addvec");
//    nastavit    parametry   thread    bloku cuFuncSetBlockShape(myKern,    x,    y,    z);
//    nakopírovat    size    bajtů    z  &ptr   do    prostoru    parametrů //    kernelu    myKern   na    pozici    offset cuParamSetv(myKern,    offset,   &ptr,    size);
//   celková    velikost    parametrů cuParamSetSize(myKern,    offset);
//   execute    kernel
cuLaunchGrid(myKern,    grid_width,    grid_height);
■I
i
•f) <\(y
Jak programovat CUDA ooooooooooo»ooooooooooooo
Specifické výhody obou rozhraní
Aneb, které rozhraní použít. Runtime API:
•  CUFFT, CUBLAS, CUDPP knihovny
•  Emulace karty Driver API:
•  Správa kontextů
•  Podpora 16-bitových float textur
•  Just-in-time (JIT) kompilace PTX kernelů
•  Přístup k MCL knihovně
□       g
Jak programovat CUDA oooooooooooo«oooooooooooo
Jak pracovat s kartami - základní funkce
Základní funkce pro výběr karty, na které bude proveden výpočet
•  cudaGetDeviceCountf/nř *count) - počet dostupných karet s compute capability > 1.0, pokud v systému není dostupná žádná karta, vrátí funkce hodnotu 1, protože systém podporuje emulační mód - compute capability bude Major: 9999 Minor: 9999
•  cudaSetDevice(7nŕ dev) - musí být voláno před inicializací, v opačném případě vrací funkce chybové hlášeni cudaErrorSetOnActiveProcess
•  cudaGetDevicef/nř *dev) - právě používané zařízení
Jak programovat CUDA ooooooooooooo»ooooooooooo
Jak pracovat s kartami - pokročilé funkce
cudaGetDeviceProperties(sřmcř cudaDeviceProp *p, int dev) - ve struktuře cudaDeviceProp vrací informace o zařízení dev
cudaChooseDevicef/nř *dev, const struct cudaDeviceProp *p) - funkce vybere kartu na základě kriterií *p
cudaSetValidDevicesf/nř *dev_arr,int len) - seznam karet, ze kterých může být vybíráno
cudaSetDeviceFlagsf/nř flags) - nastavuje jak bude CPU vlákno čekat na kartu (Spin, Yield, Sync, Auto) nebo příznak umožňující mapovat paměť. Funkce musí být volána před inicializací
s
■O Q-C^
•  Alokace paměti na kartě - cudaMalloc{Pitch, Array, 3D, 3DArray}()
•  Lineární paměť
•  2D paměť a 2D pole
•  3D paměť a 3D pole
•  Kopírování paměti mezi počítačem a kartou (host <ř» device) kopírování dat na kartě (device <ř» device) - cudaMemcpy*()
•  Alokace paměti v RAM počítače
•   K čemu?
Jak programovat CUDA ooooooooooooooo»ooooooooo
Kopírování paměti mezi počítačem a kartou
•  Základní funkce cudaMemcpy(Vo/c/ *dst, const void *src, size-t count, en u m čuda MemcpyKind kind)
9 cudaMemcpyHostToDevice
•   cudaMemcpyDeviceToHost
•   cudaMemcpyDeviceToDevice, cudaMemcpyHostToHost
•  Teoretická přenosová rychlost dosažitelná na PCI Express 2 x 16 sběrnici je 8 GB/s. Prakticky však mnohem méně.
s
Jak programovat CUDA oooooooooooooooo»oooooooo
Kopírováni dat do karty
Dva přístupy, jeden výrazně rychlejší.
int   *hmem,    * drnem;
hmem  =   (int    *)malloc(SIZE);
cudaMalloc((void**)&dmem, SIZE]
cudaMemcpy(drnem, hmem, SIZE,
cudaMemcpyHostToDevice]
int *hmem, »drnem; cudaMallocHost (( void**)&hmem ,
SIZE); cudaMalloc((void**)&dmem, SIZE]
cudaMemcpy(drnem, hmem, SIZE,
cudaMemcpyHostToDevice]
■I
S
•f)<\(y
Jak programovat CUDA ooooooooooooooooc
Kopírováni dat do karty
Dva přístupy, jeden výrazně rychlejší.
int   *hmem,    * drnem;
hmem  =   (int    *)malloc(SIZE);
cudaMalloc((void**)&dmem, SIZE]
cudaMemcpy(drnem, hmem, SIZE,
cudaMemcpyHostToDevice]
int *hmem, »drnem; cudaMallocHost (( void**)&hmem ,
SIZE); cudaMalloc((void**)&dmem, SIZE]
cudaMemcpy(drnem, hmem, SIZE,
cudaMemcpyHostToDevice]
•  PCI-e 1.0 xl6 1,5 GB/s
•  PCI-e 2.0 xl6 4,7GB/s
•  PCI-e 1.0 xl6 2,8GB/s
•  PCI-e 2.0 xl6 5,5GB/s
S
•f)<\(y
Jak programovat CUDA
Page-locked memory
Page-locked (pinned) paměť umožňuje alokovat funkce cudaMallocHost(Vo/c/ **ptr, sizeJ. size) nebo:
cudaHostAlloc(Vo/c/ **ptr, sizeJ. size, usignedt int flags) • cudaHostAllocDefault, cudaHostAllocPortable,
cudaHostAllocMapped, cudaHostAllocWriteCombined
Paměť je alokována jako souvislý blok ve fyzickém adresním
prostoru který je navíc uzamčen proti přesunu do swapovacího
oddílu
CUDA totiž může použít pouze DMA přístup, pro který je
právě potřeba, aby daný paměťový blok byl umístěn v RAM
CUDA nepodporuje ani scatter-gather DMA, kdy je možno
najednou přistoupit ke množině adres (bloků)
Toho nelze docílit kombinací volání mallocO a mlockO
(zejména souvislost nelze zajistit z US)
•f) <\(y
Jak programovat CUDA
Page-locked memory
Není-li paměť alokována tímto způsobem, musí pak driver při kopírování do karty nejprve interně přenést data do "vhodné" paměťové oblasti a odtud je teprve kopírovat do karty (pomoci DMA)
cudaHostAlloc(9 tedy:
•  Alokuje souvislý blok paměti ve fyzickém adresním prostoru (a namapuje jej do virtuální paměti aplikace)
•   Znemožní přesun této paměti do swapovací oblasti
•   Driver si navíc pro daný kontext (nebo pro všechy) pamatuje že k dané paměti lze přistoupit přímo pomoci DMA
s
■O Q-C*
Jak programovat CUDA ooooooooooooooooooo»ooooo
Souběžný běh výpočtu na GPU a CPU
Aby CPU vlákno mohlo během GPU výpočtu vykonávat další operace a nemusel vždy čekat na GPU, jsou některé CUDA funkce asynchronní. Příklad: Příprava dalších dat, zatímco probíhá výpočet nad předchozími daty. Asynchronní je:
•  Vykonání kernelu
•   Funkce s příponou Async určené ke kopírování paměti
•  Funkce vykonávající device<^device paměťové kopie
•  Funkce nastavující paměť
Jak programovat CUDA
OOOOOOOOOOOOOOOOOOOO-
Vykonání CPU funkce během GPU výpočtu
Příklad:
cudaMemcpyAsync(dev, hst, cudaMemcpyHostToDevice, 0)
cpuFunkce ();
kernelFunkce<«grid ,    block>»(dev );
cpuFunkce ();
EBB
S
•f) <\(y
Jak programovat CUDA ooooooooooooooooooooo»ooo
Překrývání GPU výpočtu a datových přenosů - použití
streams
Má-li GPU schopnost deviceOverlap je možné kopírovat z/do karty a zároveň provádět na kartě výpočet.
a Paměť musí být page-locked (pinned) • Použití streams
•   Representuje posloupnost CUDA volání
•  Volání příslušná různým streamům mohou být vykonána souběžně
•  Streamy lze synchronizovat, případně se dotazovat na stav výpočtu ve stream u
EBB
S
•f)<\(y
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]]
■I
S
•f) <\(y
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]) ;
float *   hostPtr;
cudaMallocHost((void**)&hostPtr,    2   *   size^
■I
S
•f)<\(y
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]) ;
float *   hostPtr;
cudaMallocHost((void**)&hostPtr,    2   *   size^
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, size, cudaMemcpyHostToDevice, stream[i
hostPtr
i * size
EBB
S
•f)<\(y
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]) ;
float *   hostPtr;
cudaMallocHost((void**)&hostPtr,    2   *   size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for    (int   i  =   0;    i  <   2;   ++i)
myKernel<<<100,   512,    0,    stream [i]»>
(outputDevPtr   +  i   *   size,    inputDevPtr   +  i   *   size,    size);
■I
S
•f)<\(y
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]) ;
float *   hostPtr;
cudaMallocHost((void**)&hostPtr,    2   *   size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for    (int   i  =   0;    i  <   2;   ++i)
myKernel<<<100,   512,    0,    stream [i]»>
(outputDevPtr   +  i   *   size,    inputDevPtr   +  i   *   size,    size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost , stream[i]);
□ ► 4 s  ► <■=► <  ■= ►
■O Q-C^
■I
ni
Jak programovat CUDA oooooooooooooooooooooo
Příklad překrývání GPU výpočtu a datových přenosů
cudaStream_t   stream [2];
for    (int   i  =   0;    i  <   2;   ++i)
cudaStreamCreate(&stream[i]) ;
float *   hostPtr;
cudaMallocHost((void**)&hostPtr,    2   *   size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for    (int   i  =   0;    i  <   2;   ++i)
myKernel<<<100,   512,    0,    stream [i]»>
(outputDevPtr   +  i   *   size,    inputDevPtr   +  i   *   size,    size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost , stream[i]);
cudaThreadSynchronize();
■I
i
Jak programovat CUDA
Detekce chyb
•  Všechny runtime funkce (cuda*(9) vracejí chybový kód typu cudaError_t
•  CUDA runtime udržuje pro každé CPU vlákno chybovou proměnou, která je v případě chyby přepsána chybovou hodnotou posledního volání
•   Funkce cudaGetLastError(9 vrací obsah chybové proměnné a zároveň nastaví její hodnotu na cudaSuccess
•  Chybový kód lze do slovní podoby přeložit voláním cudaGetErrorString(9
•   Návratová hodnota asynchronních funkcí lze spolehlivě ověřit pouze explicitním voláním cudaThreadSynchronizefJ a ověřením jeho návratové hodnoty
s
■O Q-C^
Jak programovat CUDA
OOOOOOOOOOOOOOOOOOOOOOOOl
Přiklad detekce chyb
cudaError_t err = cudaSetDevice ( . .. if(err != cudaSuccess) {
fprintf(stderr, "Error: '%s'\n"
exit(CHYBA);
}
//< synchronní volání cudaGetErrorString(err));
□ ► 4  rfP ► <-=► < -= ►
•f) <\(y
■I
Jak programovat CUDA
OOOOOOOOOOOOOOOOOOOOOOOOl
Přiklad detekce chyb
cudaError_t err = cudaSetDevice ( . ..) ;   //< synchronní volání if(err != cudaSuccess) {
fprintf(stderr, "Error: '%s'\n", cudaGetErrorString(err))
exit(CHYBA);
}
cudaError_t err;
cudaMemcpyAsync ( . . . ) ;   //< asynchronní volání
err = cudaThreadSynchronize();
if(err != cudaSuccess) {
fprintf(stderr, "Error: '%s'\n", cudaGetErrorString(err))
exit(CHYBA); }
■I
S
•f)<\(y
Kompilátor NVCC
•  Kompiluje CUDA zdrojové kódy obsahující CPU i GPU kód (host/device code)
•  CPU kód je předán externímu kompilátoru - gcc na linuxu, cl ve windows
•  GPU kód převeden do PTX formy, dál do binární cubin podoby
•  Výsledek GPU kompilace - .cubin výstup -je zabudován do zbytku programu
a Může být načten za běhu - viz driver API
Jak programovat CUDA
OOOOOOOOOOOOOOOOOí":
Kroky nvcc kompilace
•  Jednotlivé kroky nvcc kompilátoru lze prohlédnout, je-li kompilace spuštěna s parametry --dryrun a --keep
•  Vyzkoušej!
■I
S
•f) <\(y
Jak programovat CUDA ooooooooooooooooooooooooo
Užitečné parametry nvcc kompilátoru
•  -arch=sm_13 - zapne podporu double precission
•  -use_fast_math - automatické použití
"rychlých"matematických funkcí prefixovaných __ (vyšší rychlost nižší přesnost)
•  --ptxas-options=-v - mj. zobrazí využití registrů a paměti
•  -G - zapne debuging pro GPU kód
•  --maxrregcount < N > - nastaví maximální počet registrů, pro GPU funkce
•  -deviceemu - emulace
s
Jak programovat CUDA oooooooooooooooooooooc
Debuging CUDA aplikací
•  Obtížnější než na CPU
•  Na GPU nelze použít printf
•  Lze kopírovat mezivýsledky do globální paměti a zpět do RAM počítače - obtížné
•  Hledání chybové řádky půlením intervalů (zakomentováni řádků)
•  Emulace běhu na CPU
■I
S
•f) <\(y
OOOOOOOOOOOOOOOOOí":
Emulace běhu na CPU
•  Použitím přepínače nvcc kompilátoru lze emulovat všechen GPU kód na CPU
•  -deviceemu, --device-emulation
•  Lze použít standardní techniky hledání chyb gdb, printf
•  Vlákna jsou však spouštěna sekvenčně - nemusí se projevit chyby paralelního přístupu do paměti, odlišné paměťové modely,
•  Práce s CPU ukazatelem v GPU kódu se nemusí při emulaci projevit, avšak při běhu na GPU způsobí chybu
•  Výsledky operací v plovoucí desetinné mohou být jemně odlišné
•  Velikost warpu je rovna jedné
s
Jak programovat CUDA
•  Umožňuje za hledání chyb v aplikaci za běhu na GPU
•  Port GNU GDB 6.6
•  Velmi podobný přístup
•  Podporováno na všech kartách s compute capability 1.1 a vyšší
•  Napríklad 8800 Ultra/GTXje pouze 1.0
•  Součást CUDA Toolkit 2.3
EBB
S
•f)<\(y
ČUDA gdb
Zastavení běhu na libovolné CPU i GPU funkci nebo řádku zdrojového kódu
• (cuda-gdb) break mujKernel
• (cuda-gdb) break mujKod.cu:45
Krokování GPU kódu po warpech
•   (cuda-gdb)   next - posun po řádcích, nevkročí do funkce
•   (cuda-gdb)   step - krok do funkce
Prohlížení paměti, registrů a speciálních proměnných
•   (cuda-gdb)   print blockldx
$   l={x=0,   y=0>
s
■O Q-C*
Jak programovat CUDA ooooooooooooooooooooooooo
•  Výpis informací o použité kartě, paměti alokované na karate
•   (cuda-gdb)   info  cuda state
•  Výpis informací o blocích a vláknech běžících na kartě
•   (cuda-gdb)   info  cuda threads
•  Přepnutí na konkrétní blok nebo vlákno
•   (cuda-gdb)   thread<«Bx,   By,  Tx,  Ty,  Tz>»
s
Jak programovat CUDA
• Program musí být zkompilován s parametry -g -G
nvcc -g -G -o program program.cu
□ ► 4 s ► <■=► < ■= ►
EBB
ni
Jak programovat CUDA
CUDA Profiler
•  Umožňuje analyzovat HW čítače a odhalit neoptimální sekce kódu
•   Pro funkci umí zobrazit:
•  Čas strávený na CPU a GPU
•  Obsazení GPU
•   Počet ne/sdružených čtení/zápisů do globální paměti
•   Počet čtení/zápisů do lokální paměti
•   Počet divergentních větvení uvnitř warpu
Hodnoty jsou však měřeny pouze na jednom m u Iti procesoru, tzn. spise pro relativní porovnání mezi jednotlivými verzemi kernelu
s
oooooooooooooooooooooc
Knihovny využívající CUDA
•  Součástí CUDA instalace
•   CUBLAS - Basic Linear Algebra Subprograms (BLAS)
•   CUFFT - Fast Fourier Transform (FFT)
•  CUDPP - Data Parallel Primitives (DPP)
a http://gpgpu.org/developer/cudpp
•  Například:
•   Paralelní třídění
•   Paralelní redukce
a  Pseudonáhodný generátor čísel
•   BSD licence
s
•   Implementace BLAS pro CUDA
•  Není potřeba přímá interakce s CUDA API
•  Funkce definovány v cublas.h
•  Jednoduché použití
CUBLAS inicializace
Alokace paměti na GPU použitím CUBLAS
Naplnění alokované paměti (kopírování dat)
a
o o
•  Volání CUBLAS funkcí
•  Získání výsledků (kopírování z karty)
Ukončení CUBLAS
o simpleCUBLAS příklad v CUDA SDK
Jak programovat CUDA
•  Implementace FFT pro CUDA
•  Vyžaduje použití základních runtime API volání
(cudaMallocO ,   cudaMemcpyO)
•   Funkce definovány v cufft.h
•  ID, 2D, 3D transformace na reálných i komplexních číslech
•  simpleCUFFT příklad v CUDA SDK
■I
S
•f) <\(y
Dnes jsme si ukázali
•  Jak programovat CUDA aplikace - dvě rozhraní a rozdíly mezi nimi
•  Základní funkce runtime API
•  Jak efektivně využít šířku PCIe sběrnice při kopírování dat
•  Jak souběžně vykonávat CPU a GPU kód (překrývání)
•  Jak hledat chyby - emulace a cuda-gdb
•  Knihovny používající CUDA
Jak programovat CUDA
Samostatná práce
K samostatné práci
•  Zkuste změřit jaké rychlosti jste schopni dosáhnout při přenosu dat po PCIe sběrnici na vašem systému
•  Zkuste si vytvořit jednoduchý program, který vypíše základní informace o vaši kartě (zkuste takovýto program spustit na systému bez CUDA enabled karty)
•  Na kódu z minulé přednášky vyzkoušejte použití cuda-gdb a cudaprof