Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny OOO Závět CUDA nástroje a knihovny Jiří Matela podzim 2011 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC •OO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Rekapitulace • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Rekapitulace o»o Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny ooo Závet 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 ofthat kind is transformative. It would completely transform adjacent industries. — Jen-Hsun Huang, nVidia CEO Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Rekapitulace 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 ofthat kind is transformative. It would completely transform adjacent industries. — Jen-Hsun Huang, nVidia CEO Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny ooo Závět Rekapitulace • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OO* OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny ooo Závět Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Architektura CUDA c OpenCL Fortran C++ DX11 Compute ■ CUDA Architecture Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OO* OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Hm* U, 1) Thread {0,0) Thread (1, 0) 1 Thread (2,0) Thread (3,0) Thread {0,1) Thread (1,1) 1 Thread (2,1) Thread (3,1) I Thread (0, 2) Thread (1, 2) 1 Thread (2, 2) Thread (3, 2) 1 Jiří Matela CUDA nástroje a knihovny Hierarchie vláken Grid Block (0,0) Block (1,0) Block (2,0) Rekapitulace Jak programovat CUDA OO* OOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC Knihovny OOOOOOO OOO Závěr Rekapitulace Hierarchie pamětí • Proč programovat GPU • GPU architektura (vs. CPU) J • CUDA (Compute Unified Device Architecture) Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny ooo Závět Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) • Dvě API Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OO* OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) • Dvě API • Ukázkový příklad • Syntaktická rozšíření jazyka C - např.: __device__ • Volání runtime API -např.: cudaMallocO Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC ♦OOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Rekapitulace Jak programovat CUDA NVCC Knihovny ooo o«ooooooooooooooooooooooooo ooooooo ooo 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 kernely (moduly) • Konfigurace výpočtu • Paměťové operace • Práce s texturami • Spolupráce s OpenGL a Direct3D Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOÄOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API Rekapitulace OOO • Runtime API a C for CUDA - množina rozšíření jazyka C • Automatická inicializace, práce s kontextem a práce s kernely (moduly) • 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooo«ooooooooooooooooooooooo OOOOOOO Knihovny OOO Závět Příklad kódu používajícího CUDA rozšíření jazyka C Konfigurace CUDA kernelu addvecO int main() { addvec«(d_a , d_b , d_c ) ; } Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooo«ooooooooooooooooooooooo ooooooo Knihovny OOO Závět Příklad kódu používajícího CUDA rozšíření jazyka C Konfigurace CUDA kernelu addvecO int main() { addvec«(d_a , d_b , d_c ) ; } Překlad: $ nvcc -I/usr/local/cuda/include -L/usr/local/cuda/lib \ -lcudart -o vecadd vecadd.cu Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC ooo oooo«oooooooooooooooooooooo OOOOOOO Knihovny OOO Závěr Runtime API Knihovna runtime API: • cuda_runtime_api.h - nízkoúrovňové C funkce • cuda_runtime.h - vysoko úrovňové C++ funkce - obaluje C api Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo«oooooooooooooooooooooo OOOOOOO OOO Runtime API Rekapitulace OOO 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 Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOÄOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API Rekapitulace OOO Knihovna runtime API: • cuda_runtime_api.h - nízkoúrovňové C funkce • cuda_runtime.h - vysoko ú rov nové C++ funkce - obaluje C api • Funkce pro alokaci a dealokaci paměti • Správa karet - výběr a konfigurace karty Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOÄOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API Rekapitulace OOO Knihovna runtime API: • cuda_runtime_api.h - nízkoúrovňové C funkce • cuda_runtime.h - vysoko ú rov nové C++ funkce - obaluje C api • Funkce pro alokaci a dealokaci paměti • Správa karet - výběr a konfigurace karty • Přenos dat z/do karty Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOÄOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API Rekapitulace OOO Knihovna runtime API: • cuda_runtime_api.h - nízkoúrovňové C funkce • cuda_runtime.h - vysoko ú rov nové C++ funkce - obaluje C api • Funkce pro alokaci a dealokaci paměti • Správa karet - výběr a konfigurace karty • Přenos dat z/do karty • Debuging • Volání prefixované cuda* Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooo«ooooooooooooooooooooo ooooooo Knihovny OOO Závět Příklad kódu používajícího runtime API volání Informace o kartě int main() { cudaGetDeviceCount (&devCount ); printf(" Available devices: %d \n" , devCount ); cudaGetDeviceProperties(devProp, 0); pr intf (" Device : %d\n" , i ); printf(" Name: %s\n" , devProp—>name ); } Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooo«ooooooooooooooooooooo ooooooo Knihovny OOO Závěr Příklad kódu používajícího runtime API volání Informace o kartě int main() { cudaGetDeviceCount (&devCount ); printf(" Available devices: %d \n" , devCount ); cudaGetDeviceProperties(devProp, 0); pr intf (" Device : %d\n" , i ); printf(" Name: %s\n" , devProp—>name); } Překlad: $ gcc -I/usr/local/cuda/include -L/usr/local/cuda/lib \ -lcudart -x c -o info info.cu Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO 000000*00000000000000000000 ooooooo OOO Driver API 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 • Volání prefixované cu* Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooo«ooooooooooooooooooo ooooooo Knihovny OOO Závět Context Prostředí CUDA výpočtu představuje context • 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 (v Runtime API je však kontext svázán s CPU vláknem) • Použití více karet jedním CPU vláknem Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooo»oooooooooooooooooo ooooooo ooo 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 (Sicont , 0); // vyber první kartu (0) cuCtxCreate(&cont, CU_CTX_SCHED_AUTO, dev)); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooo«ooooooooooooooooo ooooooo Knihovny OOO Závěr Moduly 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) Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooo»oooooooooooooooo ooooooo Knihovny OOO Závět Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad (&myModule , "vectorAdd . cubin" ); cuModuleGetFunctionf&myKern , myModule , " addvec" ) ; Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooo»oooooooooooooooo ooooooo Knihovny OOO Závět 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); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooo»oooooooooooooooo ooooooo Knihovny OOO Závět Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad (&myModule, "vectorAdd.cubin"); cuModuleGetFuncti on (&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); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooo»oooooooooooooooo ooooooo Knihovny OOO Závět Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad (&myModule, "vectorAdd.cubin"); cuModuleGetFuncti on (&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); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooo»oooooooooooooooo ooooooo ooo Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad (&myModule, "vectorAdd.cubin"); cuModuleGetFuncti on (&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); 13 ► i -00.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooo«ooooooooooooooo ooooooo ooo Specifické výhody obou rozhraní Aneb, které rozhraní použít. Runtime API: • Jednodužší • CUFFT, CUBLAS, CUDPP knihovny • Emulace karty Driver API: • Správa kontextů • Větší kontrola CUDA prostředí Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooo«ooooooooooooooo ooooooo Knihovny OOO Závět Specifické výhody obou rozhraní Aneb, které rozhraní použít. Runtime API: • Jednodužší • CUFFT, CUBLAS, CUDPP knihovny • Emulace karty Driver API: • Správa kontextů • Větší kontrola CUDA prostředí Rozhraní lze kombinovat! Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooooo»oooooooooooooo ooooooo Knihovny OOO Závět Jak pracovat s kartami - základní funkce Základní funkce pro výběr karty • cudaGetDeviceCount(7nr *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 M i nor: 9999 • cudaSetDevice(7nr dev) - musí být voláno před inicializací, v opačném případě vrací funkce chybové hlášeni cudaErrorSetOnActiveProcess • cudaGetDevice(7nr *dev) - právě používané zařízení n > < & * 4 = > < = > s -00.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooo«ooooooooooooo ooooooo ooo Jak pracovat s kartami - pokročilé funkce • cudaGetDeviceProperties(sŕrucŕ cudaDeviceProp *p, int dev) - ve struktuře cudaDeviceProp vrací informace o zařízení dev • cudaChooseDevice(7nr *dev, const struct cudaDeviceProp *p) - funkce vybere kartu na základě kriterií *p • cudaSetValidDevices(7nr *dev_arr,int len) - seznam karet, ze kterých může být vybíráno • cudaSetDeviceFlags(7nr flags) - nastavuje jak bude CPU vlákno čekat na kartu (Spin, Yield, Syne, Auto) nebo příznak umožňující mapovat paměť. Funkce musí být volána před inicializací Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOOOOOOOOOOO^OOOOOOOOOOOO OOOOOOO OOO Práce s pamětí • Alokace paměti na kartě - cudaMallocjPitch, Array, 3D, 3DArray}() • Lineární pamětí • 2D pamětí a 2D pole • 3D pamětí a 3D pole • Kopírování paměti mezi počítačem a kartou (host 44> device) kopírování dat na kartě (device 44> device) - cudaMemcpy*() • Alokace paměti v RAM počítače • K čemu? Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO 000000000000000*00000000000 ooooooo ooo 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 um cudaMemcpyKind kind) o cudaMemcpyHostToDevice • cudaMemcpyDeviceToHost • cudaMemcpyDeviceToDevice, cudaMemcpyHostToHost • Teoretická přenosová rychlost dosažitelná na PCI Express 2.0 xl6 sběrnici je 8 GB/s. Prakticky však mnohem méně. Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooo«oooooooooo OOOOOOO OOO Kopírováni dat do karty Dva přístupy, jeden výrazně rychlejší. int *hmem , *dmem; hmem = (int *)malloc(SIZE); cudaMalloc((void**)&dmem , SIZE ) ; cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice); int *hmem, *dmem; cudaMallocHost((void**)&hmem, SIZE ) ; cudaMalloc((void**)&dmem, SIZE ) ; cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooo«oooooooooo OOOOOOO OOO Kopírováni dat do karty Dva přístupy, jeden výrazně rychlejší. int *hmem , *dmem; hmem = (int *)malloc(SIZE); cudaMalloc((void**)&dmem , SIZE ) ; cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice); • PCI-e 1.0 xl6 l,5GB/s • PCI-e 2.0 xl6 4,7GB/s int *hmem, *dmem; cudaMallocHost ((void**)&hmem , SIZE ) ; cudaMalloc((void**)&dmem, SIZE ) ; cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice); • PCI-e 1.0 xl6 2,8GB/s • PCI-e 2.0 xl6 5,5GB/s Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC OOOOOOOOOOOOOOOOOÍOOOOOOOOO ooooooo Knihovny OOO Závěr Page-locked memory • Page-locked (pinned) paměť umožňuje alokovat funkce cudaMallocHost(Vo/c/ **ptr, s/ze_ř size) nebo: • cudaHostAlloc(Vo/c/ **ptr, s/ze_ř 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 mlock() (zejména souvislost nelze zajistit z US) Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooooo«oooooooo ooooooo ooo 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) • cudaHostAllocfJ 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 Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závět 0000000000000000000*0000000 OOOOOOO OOO Další alokace Rekapitulace OOO • Portable memory • page-locked v kontextu všech karet • cudaHostAllocPortable flag Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závět 0000000000000000000*0000000 OOOOOOO OOO Další alokace Rekapitulace OOO • Portable memory • page-locked v kontextu všech karet • cudaHostAllocPortable flag • Write-Combining Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OOO 0000000000000000000*0000000 ooooooo Knihovny OOO Závěr Další alokace • Portable memory • page-locked v kontextu všech karet • cudaHostAllocPortable flag • Write-Combining • Mapped Memory Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO 00000000000000000000*000000 OOOOOOO OOO 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í device44>device paměťové kopie • Funkce vykonávající host44>device paměťové kopie nad daty < 64KB • Funkce nastavující paměť Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooo»ooooo OOOOOOO OOO Vykonání CPU funkce během GPU výpočtu Příklad: cudaMemcpyAsync ( dev , hst , cudaMemcpyHostToDevi.ee cpuFunkce(); kernelFunkce<«grid , block»>(dev ) ; cpuFunkce(); o); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOO^OOOO OOOOOOO 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. • 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 streamu Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo«ooo ooooooo Knihovny OOO Závěr Příklad překrývání GPU výpočtu a datových přenosů cudaStream_t streai[2]; for (int i = 0; i < 2; ++i) cudaStreamCreate(&stream[i]); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo»ooo ooooooo Knihovny OOO Závět 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; -f+i) cudaStreamCreate(&stream[i]); float * hostPtr; cudaMallocHost((void**)&hostPtr, 2 * size); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo»ooo ooooooo Knihovny OOO Závet 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; -f+i) cudaStreamCreate(&strearn[i]); float * hostPtr; cudaMallocHost((void**)&hostPtr, 2 * size); for (int i = 0; i < 2; -f+i) cudaMemcpyAsync(inputDevPtr + i * size, hostPtr size, cudaMemcpyHostToDevice, stream[i]); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo»ooo ooooooo Knihovny OOO Závet 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; -f+i) cudaStreamCreate(&strearn[i]); float * hostPtr; cudaMallocHost((void**)&hostPtr, 2 * size); for (int i = 0; i < 2; -f+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); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo»ooo ooooooo Knihovny OOO Závet 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(&strearn[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 ] ); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooo»ooo ooooooo Knihovny OOO Závet 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; -f+i) cudaStreamCreate(&strearn[i]); float * hostPtr; cudaMallocHost((void**)&hostPtr, 2 * size); for (int i = 0; i < 2; -f+i) cudaMemcpyAsync(inputDevPtr + i * size, hostPtr size , cudaMemcpyHostToDevice , stream[i ]); i * size. 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(); m -00,0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO oooooooooooooooooooooooo«oo ooooooo OOO Multi GPU • Jedna aplikace může použít více GPU • cudaSetDevicefJ Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOOOOOOOOOOOOOOOOOOOOOÄOO OOOOOOO OOO Multi GPU Rekapitulace OOO • Jedna aplikace může použít více GPU • cudaSetDevicefJ • Peer-to-Peer Memory Access • 64-bit aplikace • Compute cap. 2.x na Tesla kartách • Win Vista a 7 (v Tesla Compute Cluster Mode), Win XP, Linux • Zároveň i unifikovaný adresní prostor (host a GPU karty) Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOOOOOOOOOOOOOOOOOOOOOÄOO OOOOOOO OOO Multi GPU Rekapitulace OOO • Jedna aplikace může použít více GPU • cudaSetDevicefJ • Peer-to-Peer Memory Access • 64-bit aplikace • Compute cap. 2.x na Tesla kartách • Win Vista a 7 (v Tesla Compute Cluster Mode), Win XP, Linux • Zároveň i unifikovaný adresní prostor (host a GPU karty) • Peer-to-Peer Memory Copy • cudaMemcpyPeer*^ Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OOO OOOOOOOOOOOOOOOOOOOOOOOOO0O OOOOOOO Knihovny OOO Závěr Detekce chyb • Všechny runtime funkce (cuda*fj) 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 cudaGetLastErrorfJ 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 cudaGetErrorStringfJ • 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOO* OOOOOOO OOO Příklad detekce chyb cudaError_t err = cudaSetDevice(...); //< synchronní volání if(err != cudaSuccess) { fprintf(stderr, "Error: '% s '\n" , cudaGetErrorString(err)); exit(CHYBA ) ; } Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OOO OOOOOOOOOOOOOOOOOOOOOOOOOO* ooooooo Knihovny OOO Závěr Příklad 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 ) ; } Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO 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 • Může být načten za běhu - viz driver API Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Kroky nvcc kompilace • Jednotlivé kroky nvcc kompilátoru lze prohlédnout, je-li kompilace spuštěna s parametry --dryrun a --keep • Vyzkoušej! Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny OOO Generování kódu pro konkrétní Compute Capability • - -gpu-architecture (-arch) • virtuální architektura compute.* PTX • - -gpu-code (-code) • generuje binárni kod pro konkrétni architekturu sm_* Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny OOO Závět Generování kódu pro konkrétní Compute Capability -gpu-architecture (-arch) • virtuální architektura compute_* PTX • neni-li zadáno -code použije se odpovidajici nastaveni z -arch • napr: nvcc -arch=sm_13 ekvivalentní k nvcc -arch=compute_13 -code=sm_13 -gpu-code (-code) • generuje binárni kod pro konkrétni architekturu sm_* Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny OOO Závět Generování kódu pro konkrétní Compute Capability • - -gpu-architecture (-arch) • virtuální architektura compute_* PTX • neni-li zadáno -code použije se odpovidajici nastaveni z -arch • napr: nvcc -arch=sm_13 ekvivalentní k nvcc -arch=compute_13 -code=sm_13 • - -gpu-code (-code) • generuje binárni kod pro konkrétni architekturu sm_* • je li parametrem i virtuální architektura, pak je přibalen i PTX • napr: -arch=compute_13 -code=compute_13,sm_13 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny OOO Závět Generování kódu pro konkrétní Compute Capability • - -gpu-architecture (-arch) • virtuální architektura compute_* PTX • neni-li zadáno -code použije se odpovidajici nastaveni z -arch • napr: nvcc -arch=sm_13 ekvivalentní k nvcc -arch=compute_13 -code=sm_13 • - -gpu-code (-code) • generuje binárni kod pro konkrétni architekturu sm_* • je li parametrem i virtuální architektura, pak je přibalen i PTX • napr: -arch=compute_13 -code=compute_13,sm_13 • compute_10, compute_ll, compute_12, compute_13, compute_20, compute_30, sm_10, sm_ll, sm_12, sm_13, sm_20, sm_21, sm_22, sm_23, sm_30 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC ooo ooooooooooooooooooooooooooo ooooooo Knihovny OOO Závěr 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 (deprecated) Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC OOOOOOOOOOOOOOOOOOOOOOOOOOO «000000 Knihovny OOO Závěr Debuging CUDA aplikací • Obtížnější než na CPU • Na GPU nelze použít printf - na sm_2.x lze • 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ání řádků) • Emulace běhu na CPU (nyní už nepodporován) Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOOOOOOOOOOOOOOOOOOOOOOOO OÄOOOOO OOO CUDA gdb Rekapitulace OOO • 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/GTX je pouze 1.0 • Součást CUDA Toolkit Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC OOOOOOOOOOOOOOOOOOOOOOOOOOO OOÍOOOO Knihovny OOO Závět CUDA 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} Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO ooooooooooooooooooooooooooo ooocooo OOO CUDA gdb • 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« Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOÍOO OOO CUDA gdb • Program musí být zkompilován s parametry -g -G nvcc -g -G -o program program.cu Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooo«o Knihovny OOO Závět 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. spíše pro relativní porovnání mezi jednotlivými verzemi kernelu Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOO« OOO CUDA Profiler • NVIDIA Parallel Nsight for Visual Studio Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO ooooooo »oo 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) • http://gpgpu.org/developer/cudpp • Například: • Paralelní třídění • Paralelní redukce • Pseudonáhodný generátor čísel • BSD licence Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny o«o Závět CUBLAS • 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 volání • Naplnění alokované paměti (kopírování dat) • Volání CUBLAS funkcí o Získání výsledků (kopírování z karty) • Ukončení CUBLAS • simpleCUBLAS příklad v CUDA SDK Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooooo ooooooo Knihovny oo* Závěr CUFFT • Implementace FFT pro CUDA • Vyžaduje použití základních runtime API volání (cudaMalloc(), 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 Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Z' v aver Rekapitulace OOO 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO 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 Jiří Matela CUDA nástroje a knihovny