Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO oooooooooooooooooooooooooooo ooooooo OOO CUDA nástroje a knihovny Jiří Matela podzim 2012 Jiří Matela CUDA nástroje a knihovny Rekapitulace •OO Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny 0«0 OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO 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 ofthat kind is transformative. It would completely transform adjacent industries. — Jen-Hsun Huang, nVidia CEO Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace 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 OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace • Proč programovat GPU • GPU architektura (vs. CPU) Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OO* OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Architektura CUDA DX11 C OpenCL Fortran C++ Compute CUDA Architecture Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OO* OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Hierarchie vláken 7h,«! (0.0) Ihr Ttaad.I.O. i Eli T.,,^.3,0. lilii i Threod (1,1) IBll ISlÉ Bílili Thread (0. 2f 111111 Thread (1. 2) 11111 i Thread (2, 2) Thref.il {3, 2) 111« Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OO* OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Hierarchie pamětí Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace • Proč programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) • Dvě API Jiří Matela CUDA nástroje a knihovny Rekapitulace OO* Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr Rekapitu lace • 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 OOO Jak programovat CUDA o«oooooooooooooooooooooooooo NVCC ooooooo Knihovny OOO Závěr 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 Rekapitulace Jak programovat CUDA OOO oo»ooooooooooooooooooooooooo NVCC ooooooo Knihovny OOO Závěr Runtime API • 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 Jak programovat CUDA NVCC Knihovny OOO OOO^OOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO 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 Jak programovat CUDA NVCC Knihovny OOO OOO^OOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO 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 OOO oooo»ooooooooooooooooooooooo NVCC 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 Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOO^OOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO 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 • Funkce pro alokaci a dealokaci paměti Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOÄOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API 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 Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOÄOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API 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 Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOÄOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Runtime API 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 Jak programovat CUDA NVCC Knihovny ooo ooooo«oooooooooooooooooooooo ooooooo ooo 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 Jak programovat CUDA NVCC Knihovny ooo ooooo«oooooooooooooooooooooo ooooooo ooo 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 OOO OOOOOO0OOOOOOOOOOOOOOOOOOOOO NVCC ooooooo Knihovny OOO Závěr 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 ooooooo«oooooooooooooooooooo NVCC ooooooo Knihovny OOO Závěr 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»ooooooooooooooooooo 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 ooooooooo«oooooooooooooooooo NVCC 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 Jak programovat CUDA NVCC Knihovny ooo oooooooooo«ooooooooooooooooo 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"); cuModuleGetFunctionf&myKern , myModule , " addvec" ) ; Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooo«ooooooooooooooooo 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"); // inicializace parametrů , kopírování paměti // pole argumentů void* args [] = { &d_A , &d_B , &d_C , &N }; Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooo«ooooooooooooooooo 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"); // inicializace parametrů , kopírování paměti // pole argumentů void* args [] = { &d_A , &d_B , &d_C , &N }; // execute kernel cuLaunchKernel(vecAdd, gr blckDimX, shrMemSiz, idDimX , girdDimY , gridDimZ , blckDimY , blckDimZ , stream, args, extra); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA ooo ooooooooooo«oooooooooooooooo NVCC ooooooo Knihovny OOO Závěr Specifické výhody obou rozhraní Aneb, které rozhraní použít. Runtime API: • Jednodušší • CUFFT, CUBLAS, CUDPP knihovny • Emulace karty (odstraněno od verze 3.x) Driver API: • Správa kontextů • Větší kontrola CUDA prostředí Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA ooo ooooooooooo«oooooooooooooooo NVCC ooooooo Knihovny OOO Závěr Specifické výhody obou rozhraní Aneb, které rozhraní použít. Runtime API: • Jednodušší • CUFFT, CUBLAS, CUDPP knihovny • Emulace karty (odstraněno od verze 3.x) Driver API: • Správa kontextů • Větší kontrola CUDA prostředí Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC ooo oooooooooooo«ooooooooooooooo ooooooo Knihovny OOO Závěr 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í Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC OOO OOOOOOOOOOOOOÍOOOOOOOOOOOOOO ooooooo Knihovny OOO Závěr 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 00000000000000*0000000000000 OOOOOOO OOO Práce s pamětí • Alokace paměti na kartě - cudaMalloc{Pitch, 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*000000000000 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«ooooooooooo 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«ooooooooooo 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 Jak programovat CUDA ooo ooooooooooooooooo»oooooooooo NVCC 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«ooooooooo 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 Rekapitulace Jak programovat CUDA NVCC Knihovny Závět OOO 0000000000000000000*00000000 OOOOOOO OOO Další alokace • Portable memory • page-locked v kontextu všech karet • cudaHostAllocPortable flag Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závět ooo ooooooooooooooooooo»oooooooo ooooooo ooo Další alokace • Portable memory • page-locked v kontextu všech karet • cudaHostAllocPortable flag • Write-Combining Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA OOOOOOOOOOOOOOOOOOO0OOOOOOOO NVCC ooooooo Knihovny OOO Závěr Další a lokace • 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 ooo oooooooooooooooooooo«ooooooo NVCC ooooooo Knihovny OOO Závěr 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 vykonávající host^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»oooooo ooooooo ooo Vykonání CPU funkce během GPU výpočtu Příklad: cudaMemcpyAsync(dev, hst, cudaMemcpyHostToDevice cpuFunkce(); kernelFunkce<<(dev ) ; cpuFunkce(); o); Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooooooooooooooo«ooooo ooooooo Knihovny Závěr OOO Překrývání GPU výpočtu a datových přenosů - použití streams Má-li GPU schopnost asyncEngineCount > 0 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 Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOO0OOOO NVCC Knihovny OOOOOOO 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 ] ) ; < 3 ► * ! ► < 1 ► l -00.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOO0OOOO OOOOOOO OOO 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); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooo«oooo ooooooo ooo 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 Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooo«oooo ooooooo ooo 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 Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooo«oooo ooooooo ooo 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); 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 Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooo«oooo ooooooo ooo 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 OOO OOOOOOOOOOOOOOOOOOOOOOOOÄOOO NVCC ooooooo Knihovny OOO Závěr Multi GPU • Jedna aplikace může použít více GPU • cudaSetDevicefJ Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOÄOOO OOOOOOO OOO Multi GPU • 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 Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOÄOOO OOOOOOO OOO Multi GPU • 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 Knihovny Závěr OOO 0000000000000000000000000*00 OOOOOOO OOO Multi Kernel • Od CC 2.x • Více souběžných kernelů • Vlastnost concurrentKernels • Musí být spuštěny ze stejného kontextu • Dostatek zdrojů • Max 16 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOÄO NVCC 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 OOOOOOOOOOOOOOOOOOOOOOOOOOO* NVCC 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 ) ; } g> - nastaví maximální počet registrů, pro GPU funkce • -deviceemu - emulace (deprecated) Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC •oooooo 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 Rekapitulace Jak programovat CUDA OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC OÍOOOOO Knihovny OOO Závěr CUDA gdb • 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 Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOÄOOOO OOO 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 OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC 000*000 Knihovny OOO Závěr 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 OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC oooo»oo Knihovny OOO Závěr 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 Jak programovat CUDA OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC ooooo*o Knihovny OOO Závěr 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 OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO NVCC oooooo* Knihovny OOO Závěr CUDA Profiler • NVIDIA Parallel Nsight for Visual Studio Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO oooooooooooooooooooooooooooo ooooooo »oo Knihovny využívající CUDA • Součástí CUDA instalace • CUBLAS - Basic Linear Algebra Subprograms (BLAS) o 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 Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO o«o 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 oooooooooooooooooooooooooooo NVCC 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 Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Zaver 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 u * A & > * = > 4 = * s -00.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO OOO Samostatná práce K samostatné práci • Zkuste změřit jaké rychlosti jste schopni dosáhnout při přenosu dat pjednoduchý 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