Historie GPU
I GPU = graphics processing unitI jde o akcelerátory pro algoritmy v 3D grafice a vizualizaciI mnoho z nich puvodne vzniklo pro úcely pocítacových her
I to byla casto psychologická nevýhoda GPUI typická úloha ve vizualizaci vypadá takto
I transformování miliónu polygonuI aplikování textur o velikosti mnoha MBI projekce na framebuffer
I žádná datová závislost
Historie GPUI 1970 - ANTIC in 8-bit AtariI 1980 - IBM 8514I 1993 - Nvidia Co. založenoI 1994 - 3dfx Interactive založenoI 1995 - chip NV1 od NvidiaI 1996 - 3dfx vydalo Voodoo GraphicsI 1999 - GeForce 256 by Nvidia - podpora geometrických
transformacíI 2000 - Nvidia kupuje 3dfx InteractiveI 2002 - GeForce 4 vybaveno pixel a vertex shaderyI 2006 - GeForce 8 - unifikovaná architektura (nerozlišuje
pixel a vertex shader) (Nvidia CUDA)I 2008 - GeForce 280 - podpora dvojité presnostiI 2010 - GeForce 480 (Fermi) - první GPU postavené pro
obecné výpocty - GPGPU
Výhody GPU
(Nvidia Tesla P100)I GPU je navrženo pro soucasný beh až 3800 vláken -
virtuálne až stovek tisíc vlákenI vlákna musí být nezávislá - není zaruceno, v jakém poradí
budou zpracovánaI GPU je vhodné pro kód s intenzivními výpocty a s malým
výskytem podmínekI není zde podpora spekulativního zpracováníI není zde podpora pro cacheI GPU je optimalizováno pro sekvencní prístup do pameti
- 720 GB/s
Výhody GPU
Obrázek: Zdroj Nvidia Programming Guide
Výhody GPU
Obrázek: Rust výkonu GPU a CPU
Porovnání CPU vs. GPU
Nvidia INTEL Core i7-5930GeForce 980 Six-Core
Transistors 5 200 millions ≈ 2 000 millionsClock 1.1 GHz 3.5 GHz
Threads Num. 2048 12Peak. Perf. 4.6 TFlops ≈ 0.2 TFlopsBandwidth 224 GB/s 68 GB/s
RAM 4 GB ≈ 64 GBPower 165 W 130 W
Porovnání CPU vs. GPU
Nvidia INTEL Xeon E5-2699A V4Tesla P100 22 Cores
Transistors 15.3 billions 7.2 billionsClock 1.3 GHz 3.6 GHz
Threads Num. 3584 44Peak. Perf. 10 TFlops ≈ 0.75 TFlopsBandwidth 720 GB/s 76.8 GB/s
RAM 16 GB 1.54 TBPower 300 W 145 W
GPGPU
I GPGPU = General Purpose Computing on GPU(www.gpgpu.org)
I Lengyel, J., Reichert, M., Donald, B.R. and Greenberg,D.P. Real-Time Robot Motion Planning Using RasterizingComputer Graphics Hardware. In Proceedings ofSIGGRAPH 1990, 327-335. 1990.
I 2003 - GPGPU na bežných GPUs
Nvidia CUDA
CUDA = Compute Unified Device Architecture - Nvidia 15February 2007
I výrazne zjednodušuje programování v GPGPUI zcela odstranuje nutnost pracovat s OpenGL a formulování
úloh pomocí texturI je založena na jednoduchém rozšírení jazyka C/C++I funguje jen s kartami spolecnosti Nvidia
Je velice snadné napsat kód pro CUDA ale je potreba míthluboké znalosti o GPU aby byl výsledný kód efektivní.
CUDA architektura I.Architektura Maxwell
Obrázek: Zdroj Nvidia
CUDA architektur I.
GeForce 980
I skládá se ze 4 GPC = Graphic Processing ClusterI každý GPC obsahuje 4 SMM = Streaming
MultiprocessorsI všechny GPC mezi sebou sdílejí 2MB L2 cacheI dále obsahuje 4 GB GDDR5 RAM
CUDA architektura I.Architektura Maxwell
Obrázek: Zdroj Nvidia
CUDA architektur I.
GeForce 980 – SMMKaždý SMM se skládá z:
I 32 jader/procesoru pro jednotlivá vláknaI 64 nebo 96 kB velmi rychlé sdílené pameti a 24 kB L1
cacheI sdílená pamet’ se delí do 32 moduluI LD/ST jsou jednotky pro prístup do globální pametiI SFU jsou jednotky pro výpocty složitých funkcí jako sin,cos, tan, exp
Od hardwarové architektury se odvíjí hierarchická strukturavláken:
Vlákna v CUDA
I CUDA host je CPU a operacní pamet’I CUDA device je zarízení pro paralelní zpracování až
stovek tisíc nezávislých vláken - threadsI CUDA thread je velmi jednoduchá struktura - rychle se
vytvárí a rychle se prepíná pri zpracováníI komunikace mezi výpocetními jednotkami je hlavní
problém v paralelním zpracování datI nemužeme ocekávat, že budeme schopni efektivne
synchronizovat tisíce vlákenI CUDA architektura zavádí menší skupiny vláken zvané
bloky - blocks
Bloky a gridy
I jeden blok je zpracován na jednom multiprocesoruI vlákna v jednom bloku sdílejí velmi rychlou pamet’ s
krátkou latencíI vlákna v jednom bloku mohou být synchronizovánaI v jednom bloku muže být až 1024 vláken
I multiprocesor prepíná mezi jednotlivými vláknyI tím zakrývá latence pomalé globální pametiI zpracovává vždy ta vlákna, která mají nactena potrebná
data, ostatní nacítají
Bloky vláken jsou seskupeny do gridu - grid.
Model zpracování vláken
Obrázek: Zdroj Nvidia: Getting Started with CUDA
Pamet’ový model
Obrázek: Zdroj Nvidia: Getting Started with CUDA
Pamet’ová hierarchie
Obrázek: Zdroj Nvidia: Getting Started with CUDA
Programování v CUDA
I programování v CUDA spocívá v psaní kernelu - kernelsI kód zpracovaný jedním vláknem
I kernely nepodporují rekurziI podporují vetvení kódu, ale to muže snižovat efektivituI nemohou vracet žádný výsledekI jejich parametry nemohou být referenceI podporují šablony C++I od CUDA 2.0 podporují funkci printf !!!
Programování v CUDA1 #include <cuda . h>2 #include <cuda_runt ime_api . h>3 #include <iostream >45 __global__ void cudaVectorAddi t ion ( double∗ cuda_u ,6 const double∗ cuda_v ,7 const i n t size ,8 const i n t gr i d Idx ,9 const i n t gridDim ) ;
1011 i n t main ( i n t argc , char∗ argv [ ] )12 {13 const i n t s ize ( 1 << 20 ) ;14 double ∗host_u , ∗host_v , ∗cuda_u , ∗cuda_v ;1516 /∗ ∗∗∗17 ∗ A l l o c a t i o n on the host18 ∗ /19 host_u = new double [ s i ze ] ;20 host_v = new double [ s i ze ] ;2122 /∗ ∗∗∗23 ∗ A l l o c a t i o n on the device24 ∗ /25 i f ( cudaMalloc ( ( void∗∗ ) & cuda_u , s ize ∗ sizeof ( double ) ) != cudaSuccess | |26 cudaMalloc ( ( void∗∗ ) & cuda_v , s ize ∗ sizeof ( double ) ) != cudaSuccess )27 {28 std : : ce r r << "Unable to allocate vectors on the device." << std : : endl ;29 return EXIT_FAILURE ;30 }
Programování v CUDA1 /∗ ∗∗∗2 ∗ Set t ing−up the vec to rs3 ∗ /4 for ( i n t i = 0 ; i < s ize ; i ++ )5 {6 host_u [ i ] = i ;7 host_v [ i ] = s ize − i ;8 }9 i f ( cudaMemcpy ( ( void∗ ) cuda_u , ( void∗ ) host_u ,
10 s ize ∗ sizeof ( double ) , cudaMemcpyHostToDevice ) != cudaSuccess | |11 cudaMemcpy ( ( void∗ ) cuda_v , ( void∗ ) host_v ,12 s ize ∗ sizeof ( double ) , cudaMemcpyHostToDevice ) != cudaSuccess )13 {14 std : : ce r r << "Unable to copy data from the host to the device." << std : : endl ;15 return EXIT_FAILURE ;16 }1718 /∗ ∗∗∗19 ∗ Compute the a d d i t i o n on the CPU20 ∗ /21 for ( i n t i = 0 ; i < s ize ; i ++ )22 host_u [ i ] += host_v [ i ] ;
Programování v CUDA12 /∗ ∗∗∗3 ∗ Run the CUDA kerne l4 ∗ /5 dim3 cudaBlockSize ( 256 ) ;6 cudaDeviceProp p r o p e r t i e s ;7 cudaGetDevicePropert ies ( &proper t i es , 0 ) ;8 i n t maxCudaGridSize ( p r o p e r t i e s . maxGridSize [ 0 ] ) ;9 const i n t cudaBlocksCount = s ize / cudaBlockSize . x +
10 ( s ize % cudaBlockSize . x != 0 ) ;11 const i n t cudaGridsCount = cudaBlocksCount / maxCudaGridSize +12 ( cudaBlocksCount % maxCudaGridSize != 0 ) ;13 for ( i n t g r i d I d x = 0; g r i d I d x < cudaGridsCount ; g r i d I d x ++ )14 {15 dim3 cudaGridSize ;16 i f ( g r i d I d x < cudaGridsCount )17 cudaGridSize . x = maxCudaGridSize ;18 else19 cudaGridSize . x = cudaBlocksCount % maxCudaGridSize ;20 cudaVectorAddi t ion <<< cudaGridSize , cudaBlockSize >>>21 ( cuda_u , cuda_v , s ize , g r i d Idx , maxCudaGridSize ) ;22 cudaError e r r = cudaGetLastError ( ) ;23 i f ( e r r != cudaSuccess )24 {25 std : : ce r r << "Computation on the device failed with error: "26 << cudaGetErrorSt r ing ( e r r ) << "." << std : : endl ;27 return EXIT_FAILURE ;28 }29 }
Programování v CUDA12 /∗ ∗∗∗3 ∗ Copy the r e s u l t back to the host4 ∗ /5 i f ( cudaMemcpy ( ( void∗ ) host_v , ( void∗ ) cuda_u ,6 s ize ∗ sizeof ( double ) , cudaMemcpyDeviceToHost ) != cudaSuccess )7 {8 s td : : ce r r << "Unabel to copy data back from the GPU." << std : : endl ;9 return EXIT_FAILURE ;
10 }1112 i f ( memcmp( ( void∗ ) host_u , ( void∗ ) host_v , s i ze ∗ sizeof ( double ) ) != 0 )13 {14 std : : ce r r << "The results are different." << std : : endl ;15 return EXIT_FAILURE ;16 }17 std : : cout << "Everything is ok." << std : : endl ;1819 /∗ ∗∗∗20 ∗ Freeing a l l o c a t e d memory21 ∗ /22 delete [ ] host_u ;23 delete [ ] host_v ;24 cudaFree ( cuda_u ) ;25 cudaFree ( cuda_v ) ;26 }2728 __global__ void cudaVectorAddi t ion ( double∗ cuda_u ,29 const double∗ cuda_v ,30 const i n t size ,31 const i n t gr i d Idx ,32 const i n t gridDim )33 {34 const i n t t i d = ( g r i d I d x ∗ gridDim + b lock Idx . x ) ∗ blockDim . x + th read Idx . x ;35 i f ( t i d < s ize )36 cuda_u [ t i d ] += cuda_v [ t i d ] ;37 }
Programování v CUDA
Kód uložíme do cuda-vector-addition.cu a preložímepomocí nvcc.
Vývoj efektivního kódu
Pro získání efektivního kódu je nutné dodržet následujícípravidla:
I redukovat prenos dat mezi CPU (CUDA host) a GPU(CUDA device)
I optimalizovat prístup do globální pametiI omezit divergentní vláknaI zvolit správnou velikost bloku
Komunikace mezi CPU a GPU
I komunikace pres PCI Express je velmi pomalá cca. 8 GB/sI je nutné tuto komunikaci minimalizovat
I ideálne provést jen na zacátku a na konci výpoctuI GPU se nevyplatí pro úlohy s nízkou aritmetickou
intenzitouI z tohoto pohledu mohou mít výhodu on-board GPU, které
sdílí operacní pamet’I pokud je nutné provádet casto komunikaci mezi CPU a
GPU pak je dobré jí provádet formou pipelininguI je možné provádet najednou
I výpocet na GPUI výpocet na CPUI kopírování dat z CPU do GPUI kopírování dat z GPU na CPU
Sloucené prístupy do pametí
I vetšinu prístupu GPU do globální pameti tvorí nacítánítextur
I GPU je silne optimalizováno pro sekvencní prístup doglobální pameti
I programátor by se mel vyhnout náhodným prístupum doglobální pameti
I ideální postup je:I nacíst data do sdílené pameti multiprocesoruI provést výpoctyI zapsat výsledek do globální pameti
I sloucený prístup - coalesced memory access - muževelmi výrazne snížit (až 32x) pocet pamet’ových transakcí
Sloucené prístupy do pameti
Obrázek: Zdroj Nvidia: Nvidia CUDA programming guide
Sloucené prístupy do pameti
Obrázek: Zdroj Nvidia: Nvidia CUDA programming guide
Sloucené prístupy do pameti
Obrázek: Zdroj Nvidia: Nvidia CUDA programming guide
Keš
Architektura Fermi zavádí plne funkcní L1 a L2 keše.I L1 keš se nachází na každém multiprocesoru
I lze nastavit, jaká cást ze 64kB SRAM pameti bude urcenopro keš pomocí funkce:
I cudaFuncSetCacheConfig( MyKernel,cudaFuncCachePreferShared )
I cudaFuncCachePreferShared - shared memory is 48 KBI cudaFuncCachePreferL1 - shared memory is 16 KBI cudaFuncCachePreferNone - no preference
I L2 keš je spolecná pro všechny multiprocesory a mávelikost 768kB
Sdílená pamet’ multiprocesoru
I sdílená pamet’ multiprocesoru je rozdelena na 32pamet’ových bank
I data se ukládají do jednotlivých bank vždy po 4 bajtechI je potreba se vyhnout situaci, kdy dve vlákna ze skupiny
32 ctou z ruzných adres v jedné banceI nevadí, když cte více vláken ze stejné adresy, použije se
broadcast
Divergentní vlákna
I CUDA device umí zpracovávat soucasne ruzné kernely, alejen na ruzných multiprocesorech
I Nvidia tuto architekturu nazývá SIMT = SingleInstruction, Multiple Threads
I v rámci jednoho multiprocesoru jde ale o SIMDarchitekturu, tj. všechny jednotky provádejí stejný kód
I warp je skupina 32 vláken zpracovávaných soucasneI vlákna ve warpu jsou tedy implicitne synchronizovanáI všechna by mela zpracovávat stejný kód
Zpracování bloku vláken na multiprocesoruI na mutliprocesoru vetšinou beží více bloku vlákenI scheduler mezi nimi prepíná a spouští vždy ty bloky
vláken, které mají nacteny potrebná dataI tím se zakrývají velké latence globální pameti
I k tomu je ale potreba, aby jeden blok nevycerpal všechnyregistry a sdílenou pamet’
I pokud není dostatek registru, ukládají se promenné dolocal memory - to je pomalé
I je potreba dobre zvolit velikost bloku - násobek 32I minimalizovat pocet promenných a množství sdílené pameti
použité jedním blokemI minimalizovat velikost kódu kernelu
I efektivnost obsazení multiprocesoru udává parametr zvanýoccupancy (maximum je 1.0)
I za úcelem optimalizace lze použítI CUDA occupancy calculator 1
I CUDA profilerI výpisy nvcc -ptxas-options=-v
1http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls
Multi-GPU systémy
I do jednoho PC lze instalovat az 4 GPUI navíc existují i dvoucipové GPU kartyI dostáváme tak systém s distribuovanou pametí
I jednotlivá GPU nemají prímý prístup do pameti techostatních
I více v CUDA Programming Guide