Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr ooo ooooooooooooooooooooooooo ooooooo ooo CUDA nástroje a knihovny Jiří Matela podzim 2009 □ ť5> - = = ^c^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooooooooooooooooooooo ooooooo ooo Rekapitulace Rekapitulace •oo • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Rekapitulace o*o Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooooooo Knihovny ooo Závěr 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 Jak programovat CUDA NVCC Knihovny Závěr o»o ooooooooooooooooooooooooo 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 of that kind is transformative. It would completely transform adjacent industries. — Jen-Hsun Huang, nVidia CEO □ <3 = -s -00,0 Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooooooooooooooooooooo ooooooo ooo Rekapitulace Rekapitulace oo» • Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooooooooooooooooooooo ooooooo ooo Rekapitulace Rekapitulace oo» • Proč programovat GPU • GPU architektura (vs. CPU) Jiří Matela CUDA nástroje a knihovny Rekapitulace 00» Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooooooo Knihovny ooo Závěr Rekapitulace Architektura CUDA • Proč programovat GPU o GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) c OpenCL _ Fortran C++ DX11 Compute ■ CUDA Architecture Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CL) DA NVCC ooooooooooooooooooooooooo ooooooo Knihovny ooo Závěr Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) 9 CUDA (Compute Unified Device Architecture) Boded, D Thread (0, 0) Thread (1, 0) 1 Thread (2,0) Thread (3, 0) i Thread (0,1) Thread (1,1) i Thread (2,1) Thread (3,1) 1 Thread (0, 2) Thread (l, 2) 1 Thread (2,2) Thread (3, 2) 1 ► 4 1 ► = : Jiří Matela CUDA nástroje a knihovny Hierarchie vláken Grid Block (0, 0) Block (1,0) Block (2,0) Iliil mil unii Block (0, iy' Block (1,1) -Block (2,1) Iii Jak programovat CUDA NVCC Knihovny Závěr ooooooooooooooooooooooooo ooooooo ooo Rekapitulace Rekapitulace oo« • Proc programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Hierarchie pamětí OrfdO Bock (0,0) Block (1,0) Block (2,0 Block (0,1) Block (1,1) Block (2, i: Grid 1 Global memory Block (0, 0) Block (1,0) Block (0, 1) Block (1. 1) JUHU Block (0, 2) Block (1,2) flP - "* ! -00,0 Jiří Matela CUDA nástroje a knihovny Rekapitulace oo» Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooooooo Knihovny ooo Závěr Rekapitulace • Proč programovat GPU • GPU architektura (vs. CPU) a CUDA (Compute Unified Device Architecture) • Dvě API Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooooooooooooooooooooo ooooooo ooo Rekapitulace Rekapitulace oo» • 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC •OOOOOOOOOOOOOOOOOOOOOOOO OOOOOOO Knihovny OOO Závěr Runtime API vs. Driver API Vytvářet CUDA aplikace lze užitím buďto Runtime API nebo Driver API. Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC o»ooooooooooooooooooooooo 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oo»oooooooooooooooooooooo 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooosooooooooooooooooooooo ooooooo Závěr Příklad kódu používajícího CUDA rozšíření jazyka C Konfigurace CUDA kernelu addvecO int mainQ { addvec«(d_a, d_b , d_c ); } Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC ooosooooooooooooooooooooo ooooooo Knihovny ooo Závěr Příklad kódu používajícího CUDA rozšíření jazyka C Konfigurace CUDA kernelu addvecO int mainQ { 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo»oooooooooooooooooooo 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 □ gl - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo»oooooooooooooooooooo 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo»oooooooooooooooooooo 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 • Správa karet - výběr a konfigurace karty □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo»oooooooooooooooooooo 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 • Správa karet - výběr a konfigurace karty • Přenos dat z/do karty □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr oooo»oooooooooooooooooooo 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 • Správa karet - výběr a konfigurace karty • Přenos dat z/do karty • Podpora emulace karty na CPU - debuging • Volání prefixováné cuda* □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC ooooo»ooooooooooooooooooo ooooooo Knihovny ooo Závěr Příklad kódu používajícího runtime API volání Informace o kartě int mainQ { cudaGetDeviceCount(&devCount ); pr int f ("Available devices: %d \ n " , de vCount ); cudaGetDevic ePropert ies(devProp, 0); printf (" Device : %d\n" , i ); printf("Name: %s\n" , devProp—>name ); } Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC ooooo»ooooooooooooooooooo ooooooo Knihovny ooo Závěr Příklad kódu používajícího runtime API volání Informace o kartě int mainQ { cudaGetDeviceCount (&devCount); pr int f ("Available devices: %d \ n " , de vCount ); cudaGetDevic ePropert ies(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 \ -lcudart -x c -o info info.cu □ - = = -^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooo»oooooooooooooooooo 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 • Žádná GPU emulace • Volání prefixováné cu* Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooo»ooooooooooooooooo ooooooo ooo Context Rekapitulace ooo 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 (v Runtime API je však kontext svázán s CPU vláknem) • Použití více karet jedním CPU vláknem □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooo»oooooooooooooooo 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 (&cont , 0); // vyber první kartu (0) cuCtxCreate(&cont, CU_CTX_SCHED_AUTO, dev)); □ - = = -0*3*0 Jiří Matela CUDA nástroje a knihovny Jak programovat CUDA NVCC Knihovny Závěr ooooooooo»ooooooooooooooo ooooooo ooo Moduly Rekapitulace OOO 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»oooooooooooooo ooooooo Knihovny ooo Závěr Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad(&myModul e, "vectorAdd.cubin"); cuModuleGetFunction(&myKern, myModule, "addvec"); □ - = = -0*3*0 Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooo»oooooooooooooo ooooooo Knihovny ooo Závěr Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad(&myModul e, "vectorAdd.cubin"); cuModuleGetFunction(&myKern, myModule, "addvec"); // nastavit parametry thread bloku cuFuncSetBlockShape(myKern, x, y, z); □ - = = -0*3*0 Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooo»oooooooooooooo ooooooo Knihovny ooo Závěr Příklad konfigurace a spuštění modulu CUmodule myModule; CUfunction myKern; //nahrát modul, získat kernel "addvec" cuModuleLoad(&myModul e, "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); □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooo»oooooooooooooo ooooooo Knihovny ooo Závěr 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); □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooo»oooooooooooooo ooooooo Knihovny ooo Závěr 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); □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooo»ooooooooooooo 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í □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooo»ooooooooooooo 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í Rozhraní lze kombinovat! Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooooo«oooooooooooo ooooooo Knihovny ooo Závěr Jak pracovat s kartami - základní funkce Základní funkce pro výběr karty • 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 • cudaSetDevicef/nř 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í □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC 0000000000000*00000000000 ooooooo Knihovny ooo Závěr 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 9 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, Syne, Auto) nebo příznak umožňující mapovat paměť. Funkce musí být volána před inicializací Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooo«oooooooooo ooooooo ooo Práce s pamětí • 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? □ - = = -0*3*0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooo»ooooooooo 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 u m cud a MemcpyKind kind) 9 cudaMemcpyHostToDevice • cudaMemcpyDeviceToHost • cudaMemcpyDeviceToDevice, cudaMemcpyHostToHost • Teoretická přenosová rychlost dosažitelná na PCI Express 2.0 xl6 sběrnici je 8GB/s. Prakticky však mnohem méně. □ S ~ = -š -00,0 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC oooooooooooooooo»oooooooo ooooooo Knihovny ooo Kopírováni dat do karty Dva přístupy, jeden výrazně rychlejší. int *hmem, *dmem; hmem = (int *)mal1oc(SIZE ); cudaMalloc((void**)&dmem, SIZE); cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice ); int *hmem, *dmem; cudaMallocHost((vo id* *)&hmem, SIZE ); cudaMalloc((void**)&dmem, SIZE); cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice ); □ gl - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooo»oooooooo ooooooo ooo Kopírováni dat do karty Dva přístupy, jeden výrazně rychlejší. int *hmem, *dmem; hmem = (int *)mal1oc(SIZE ); cudaMalloc((void**)&dmem, SIZE); cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice ); o PCI-e 1.0 xl6 l,5GB/s o PCI-e 2.0 xl6 4,7GB/s int *hmem, *dmem; cudaMallocHost((vo id* *)&hmem, SIZE ); cudaMalloc((void**)&dmem, SIZE); cudaMemcpy(dmem, hmem, SIZE, cudaMemcpyHostToDevice ); » PCI-e 1.0 xl6 2,8GB/s o PCI-e 2.0 xl6 5,5GB/s Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooo»ooooooo ooooooo ooo Page-locked memory • Page-locked (pinned) paměť umožňuje alokovat funkce cudaMallocHost(Vo/c/ **ptr, sizeJ. size) nebo: • cudaHostAlloc(Vo/c/ **ptr, s/ze_ř size, usignedt int fíags) • 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) Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooooosoooooo 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) • 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA ooo 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 vykonávající host<^device paměťové kopie nad daty < 64KB • Funkce nastavující paměť □ S ~ = -š -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooooooo»oooo ooooooo ooo Vykonání CPU funkce během GPU výpočtu Příklad: cudaMemcpyAsync(dev, cpuFunkce (); kernelFunkce«(dev ); o); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooo»ooo 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. 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 □ S ~ = -š -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooooooooo»oo 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 ]); □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC 0000000000000000000000*00 ooooooo Knihovny ooo Závěr 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); □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooooooooooooooo»oo ooooooo Knihovny ooo Závěr 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 ]); □ - = = -0*3*0 Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooooooooooooooo»oo ooooooo Knihovny ooo Závěr 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); □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo oooooooooooooooooooooo»oo 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(&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 ]); □ - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace ooo Jak programovat CUDA NVCC oooooooooooooooooooooo»oo ooooooo Knihovny ooo Závěr 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 ]); cudaThreadSynchronize (); □ - nastaví maximální počet registrů, pro GPU funkce • -deviceemu - emulace □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOO »000000 ooo 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ání řádků) • Emulace běhu na CPU □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooooo o»ooooo ooo 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é □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC OOOOOOOOOOOOOOOOOOOOOOOOO oo»oooo Knihovny ooo 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/GTXje pouze 1.0 • Součást CUDA Toolkit 2.3 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooo»ooo Knihovny 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 a (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 OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooo oooo«oo Knihovny 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« □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooooo»o Knihovny ooo CUDA gdb • Program musí být zkompilován s parametry -g -G nvcc -g -G -o program program.cu □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny OOO OOOOOOOOOOOOOOOOOOOOOOOOO 000000« ooo 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 □ - = = ^Q^O Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooooo 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) a http://gpgpu.org/developer/cudpp • Například: • Paralelní třídění • Paralelní redukce a Pseudonáhodný generátor čísel a BSD licence □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooooo 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í • Získání výsledků (kopírování z karty) • Ukončení CUBLAS o simpleCUBLAS příklad v CUDA SDK Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny ooo ooooooooooooooooooooooooo ooooooo oot 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 ooo ooooooooooooooooooooooooo ooooooo ooo Záver 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 □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny Rekapitulace OOO Jak programovat CUDA NVCC ooooooooooooooooooooooooo ooooooo Knihovny ooo Závěr 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 □ g - = = -0*3.0 Jiří Matela CUDA nástroje a knihovny