Rekapitulace Jak programovat C (JDA NVCC Knihovny Závěr oo oooooooooooooooooo oooooo ooo CUDA nástroje a knihovny Jiří Matela podzim 2017 Jiří Matela CUDA nástroje a knihovny Jak programovat C (JDA oooooooooooooooooo NVCC oooooo Knihovny ooo Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Jak programovat C (JDA oooooooooooooooooo NVCC oooooo Knihovny ooo Proč programovat GPU Jiří Matela CUDA nástroje a knihovny Jak programovat C (JDA oooooooooooooooooo NVCC oooooo Knihovny ooo Proč programovat GPU GPU architektura (vs. CPU) Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat C U DA oooooooooooooooooo Rekapitulace NVCC oooooo Knihovny ooo Závěr Architektura CUDA • Proc programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) 1^ c OpenCL Fortran C++ DX11 Compute ... CUDA Architecture Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat C U DA oooooooooooooooooo Rekapitulace NVCC oooooo Knihovny ooo Závěr Hierarchie vláken • Proc programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Grid Block (0,0) Block (1,0) Block (2,0) Block (0, iy' Slock (1,1) -Block (2,1) Block (1,1) Thread (0, 0) Thread (1, 0) Thread (2,0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) 1 Thread (l, 2) Thread (2, 2) Thread (3, 2) Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Rekapitulace Rekapitulace Jak programovat CUDA NVCC o« oooooooooooooooooo oooooo • Proc programovat GPU • GPU architektura (vs. CPU) • CUDA (Compute Unified Device Architecture) Hierarchie pamětí Thread Thread Block ^ GridO Block (0, 0) Block (1, 0) Block (2, 0) Block 1.0. 1) Block (1,1) Block (2,1) Grid 1 Block (0,0) Block (1,0) Block (0,1) Block (1.1} Block (0,2) Block (1,2} Per-thread local memory Per-block shared memory Global memory Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Runtime API vs. Driver API Rekapitulace Jak programovat CUDA NVCC OO «00000000000000000 oooooo Vytvářet CUDA aplikace lze užitím buďto Runtime API nebo Driver API. CPU Aplikace Aplikace 1 CUDA Runtime API r CUDA Driver API Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Možnosti rozhraní Rekapitulace Jak programovat CUDA NVCC oo o«oooooooooooooooo oooooo Rozhraní umožňují provádět na úrovni hostitelského systému (kód vykonávaný na CPU) následující operace 9 Správa zařízení • Práce s kontextem • Práce s kernely (moduly) o Konfigurace výpočtu • Paměťové operace • Práce s texturami • Spolupráce s OpenGL a Direct3D Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Runtime API Rekapitulace Jak programovat CUDA NVCC OO OO0OOOOOOOOOOOOOOO oooooo 9 Runtime API a C for CUDA - množina rozšírení jazyka C • Automatická inicializace, práce s kontextem a práce s kernely (moduly) 9 Konfigurace výpočtu (volání kernelu) - syntaktický konstrukt (rozšíření jazyka C) 9 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 Závěr OO OOO0OOOOOOOOOOOOOO oooooo 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 Závěr OO OOO0OOOOOOOOOOOOOO oooooo 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 NVCC Knihovny Závěr oo oooo«ooooooooooooo oooooo ooo Příklad kódu používajícího runtime API volání Informace o kartě int main() { cudaGetDeviceCount (&devCount); printf( " A va i I abIe devices: %d\n" , devCount); cudaGetDeviceProperties(devProp, 0); printf (" Device : %d\n" , i ) ; printf (" Name : %s\n" , devProp—>name ) ; } □ [31 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr oo oooo«ooooooooooooo oooooo ooo Příklad kódu používajícího runtime API volání Informace o kartě int main() { cudaGetDeviceCount (&devCount); printf( " A va i I abIe devices: %d\n" , devCount); cudaGetDeviceProperties(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 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OO OOOOO0OOOOOOOOOOOO oooooo ooo Jak pracovat s kartami - základní funkce Základní funkce pro výběr karty o cudaGetDeviceCount(7r?r *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(7r?r dev) - musí být voláno před inicializací, v opačném případě vrací funkce chybové hlášeni cudaErrorSetOnActiveProcess o cudaGetDevice(7r?r *dev) - právě používané zařízení Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Jak pracovat s kartami - pokročilé funkce Rekapitulace Jak programovat CUDA NVCC OO OOOOOO0OOOOOOOOOOO oooooo • cudaGetDeviceProperties(sŕ/r/cŕ cudaDeviceProp *p, int dev) - ve struktuře cudaDeviceProp vrací informace o zařízení dev • cudaChooseDevice(7r?r *dev, const struct cudaDeviceProp *p) - funkce vybere kartu na základě kriterií *p • cudaSetValidDevices(7r?r *dev_arr,int len) - seznam karet, ze kterých může být vybíráno • cudaSetDeviceFlags(7r?r flags) - nastavuje jak bude CPU vlákno čekat na kartu (Spin, Yield, BlockingSync, 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 Knihovny Závěr ooo Práce s pamětí Rekapitulace Jak programovat CUDA NVCC OO 0000000*0000000000 oooooo o Alokace paměti na kartě - cudaMalloc{Pitch, Array, 3D, 3DArray}() • Lineární paměť o 2D paměť a 2D pole o 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 44> device) - cudaMemcpy*() o Alokace paměti v RAM počítače • K čemu? Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Kopírování paměti mezi počítačem a kartou Rekapitulace Jak programovat CUDA NVCC oo oooooooo«ooooooooo oooooo • Základní funkce cudaMemcpy(Vo/c/ *dst, const void *src, size-t count, en u m cudaMemcpyKind kind) o cudaMemcpyHostToDevice o 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 Knihovny Závěr ooo Kopírováni dat do karty Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOO0OOOOOOOO oooooo 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 Knihovny Závěr ooo Kopírováni dat do karty Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOO0OOOOOOOO oooooo 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 Knihovny Závěr ooo Page-locked memory Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOOO0OOOOOOO oooooo • Page-locked (pinned) paměť umožňuje alokovat funkce cudaMallocHost(Vo/c/ **ptr, s/ze_r size) nebo: • cudaHostAlloc(Vo/c/ **ptr, s/ze_r s/ze, 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 9 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 Knihovny Závěr ooo Page-locked memory Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOOOO0OOOOOO oooooo • 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 Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Souběžný běh výpočtu na GPU a CPU Rekapitulace Jak programovat CUDA NVCC oo oooooooooooo«ooooo oooooo 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: o Vykonání kernelu • Funkce s příponou Async určené ke kopírování paměti o Funkce vykonávající device<^>device paměťové kopie o 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 Závěr OO OOOOOOOOOOOOO0OOOO oooooo ooo Vykonání CPU funkce během GPU výpočtu Příklad: cudaMemcpyAsync(dev, hst, cudaMemcpyHostToDevice , cpuFunkce () ; kernelFunkce«(dev ) ; cpuFunkce () ; □ [31 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr oo oooooooooooooo«ooo oooooo 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) 9 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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]); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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 Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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); for (int i = 0; i < 2; ++i) cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice , stream[i]); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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); 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); Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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); 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 Jak programovat CUDA NVCC Knihovny Závěr OO OOOOOOOOOOOOOOO0OO oooooo 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); 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]»> (outputDe vPt r + 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(); Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Detekce chyb Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOOOOOOOOO0O oooooo • 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 9 Chybový kód lze do slovní podoby přeložit voláním cudaGetErrorString(9 o Návratová hodnota asynchronních funkcí lze spolehlivě ověřit pouze explicitním voláním cudaThreadSynchronize(9 a ověřením jeho návratové hodnoty Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Příklad detekce chyb Rekapitulace Jak programovat CUDA NVCC OO 00000000000000000« oooooo 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 Knihovny Závěr ooo Příklad detekce chyb Rekapitulace Jak programovat CUDA NVCC OO 00000000000000000« oooooo 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 Knihovny Závěr ooo Kompilátor NVCC Rekapitulace Jak programovat CUDA NVCC oo oooooooooooooooooo oooooo o Kompiluje CUDA zdrojové kódy obsahující CPU i GPU kód (host/device code) o 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 Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr oo oooooooooooooooooo oooooo ooo Užitečné parametry nvcc kompilátoru 9 --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 Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo Debuging CUDA aplikací Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOOOOOOOOOOO «00000 o Obtížnější než na CPU 9 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é o Hledání chybové řádky půlením intervalů (zakomentování řádků) Jiří Matela CUDA nástroje a knihovny Rekapitulace oo Jak programovat CUDA oooooooooooooooooo CUDA gdb NVCC o«oooo Knihovny ooo Závěr Umožňuje za hledání chyb v aplikaci za běhu na GPU • Port GNU GDB 6.6 o Velmi podobný přístup 9 Podporováno na všech kartách s compute capability 1.1 a vy ssi • Napríklad 8800 Ultra/GTX je pouze 1.0 • Součást CUDA Toolkit Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo CUDA gdb Rekapitulace Jak programovat CUDA NVCC OO OOOOOOOOOOOOOOOOOO OO0OOO • Zastavení běhu na libovolné CPU i GPU funkci nebo řádku zdrojového kódu • (cuda-gďb) break mujKernel • (cuda-gdb) break mujKod.cu:45 o 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 $ 1 = {x = 0, y = 0} Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo CUDA gdb Rekapitulace Jak programovat CUDA NVCC oo oooooooooooooooooo ooo«oo • 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 Knihovny Závěr ooo CUDA gdb Rekapitulace Jak programovat CUDA NVCC oo oooooooooooooooooo oooo«o Program musí být zkompilován s parametry -g nvcc -g -G -o program program.cu □ |3> Jiří Matela CUDA nástroje a knihovny Knihovny Závěr ooo CUDA Profiler Rekapitulace Jak programovat CUDA NVCC oo oooooooooooooooooo ooooo* • Umožňuje analyzovat HW čítače a odhalit neoptimální sekce kódu • Pro funkci umí zobrazit: ■v • Cas 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 multiprocesoru, 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 Závěr OO OOOOOOOOOOOOOOOOOO OOOOOO «00 Knihovny využívající CUDA 9 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 o Například: • Paralelní třídění • Paralelní redukce • Pseudonáhodný generátor čísel © BSD licence Jiří Matela CUDA nástroje a knihovny Rekapitulace oo Jak programovat C (JDA oooooooooooooooooo CUBLAS NVCC oooooo Knihovny Závěr • Implementace BLAS pro CUDA • Není potřeba přímá interakce s CUDA API • Funkce definovány v cublas.h • Jednoduché použití • CUBLAS inicializace o 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 • simpleCUBLAS příklad v CUDA SDK Jiří Matela CUDA nástroje a knihovny Rekapitulace oo Jak programovat C (JDA oooooooooooooooooo CUFFT NVCC oooooo Knihovny Závěr o Implementace FFT pro CUDA o Vyžaduje použití základních runtime API volání (cudaMal1o c(), cudaMemcpy()) • Funkce definovány v cufft.h • ID, 2D, 3D transformace na reálných i komplexních číslech o simpleCUFFT příklad v CUDA SDK Jiří Matela CUDA nástroje a knihovny Rekapitulace Jak programovat CUDA NVCC Knihovny Závěr oo oooooooooooooooooo oooooo ooo Závěr Dnes jsme si ukázali 9 Jak programovat CUDA aplikace 9 Základní funkce runtime API • Jak efektivně využít šírku 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 - cuda-gdb • Knihovny používající CUDA Jiří Matela CUDA nástroje a knihovny Rekapitulace oo Jak programovat C (JDA oooooooooooooooooo Samostatná práce NVCC oooooo Knihovny ooo Závěr 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) 9 Na kódu z minulé přednášky vyzkoušejte použití cuda-gdb a cudaprof Jiří Matela CUDA nástroje a knihovny