´ Uvod
Hardware
Knihovny
NNSU
Vyuˇzit´ı GPU k urychlen´ı uˇcen´ı neuronov´ych s´ıt´ı Semin´aˇr strojov´eho uˇcen´ı a modelov´an´ı ˇ V. Spanihel Katedra matematiky Fakulta jadern´ a a fyzik´ alnˇ e inˇzen´ yrsk´ a ˇ e vysok´ Cesk´ e uˇ cen´ı technick´ e v Praze
24.10.2013
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
1 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Obsah
´ 1 Uvod 2 Hardware
3 Knihovny
4 NNSU
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
2 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Paraleln´ı programov´an´ı na GPU
NVIDIA CUDA • Pˇrehledn´ e rozhrann´ı • Pouze pro karty od firmy
NVIDIA
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
OpenCL • Hardwarovˇ e nez´avisl´e • Sloˇ zitˇejˇs´ı rozhrann´ı
3 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Co je potˇreba Software • CUDA Runtime - Nutn´ e pro spuˇstˇen´ı CUDA aplikace • CUDA Toolkit - Pro v´ yvoj(vˇsechny knihovny a dokumentace) • Podporovan´ y operaˇcn´ı syst´em s C/C++ pˇrekladaˇcem • Volitelnˇ e GPU Computing SDK - Stovky pˇr´ıklad˚ u
Hardware • Grafick´ a karta NVIDIA podporuj´ıc´ı CUDA architekturu
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
4 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Z´akladn´ı pojmy Pojmy • Host - Oznaˇ cen´ı pro CPU • Device - Oznaˇ cen´ı pro GPU • Kernel - Oznaˇ cen´ı metody spouˇstˇen´e na GPU • Vl´ akno - Jedno vol´an´ı metody • Blok - Aˇ z trojrozmˇern´a struktura vl´aken • Warp - Soubor vl´ aken zpracov´avan´ych paralelnˇe v r´amci bloku • Mˇr´ıˇ zka - Aˇz trojrozmˇern´a struktura blok˚ u
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
5 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
V´ypoˇcetn´ı schopnost (Compute capability) Vysvˇetlen´ı • Specifikace verze CUDA architektury (napˇr. CC = 2.1) • Pˇr´ıklad odliˇsnost´ı: maxim´ aln´ı poˇcet blok˚ u v gridu a vl´aken v
bloku
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
6 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
V´ypoˇcty na GPU
V´yvoj GPU architektur v ˇcase • 2008: Tesla - V´ ypoˇcetn´ı schopnost = 1.x • 2010: Fermi - V´ ypoˇcetn´ı schopnost = 2.x • 2012: Kepler - V´ ypoˇcetn´ı schopnost = 3.x • 2014: Maxwell - V´ ypoˇcetn´ı schopnost = 4.x
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
7 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
V´yvoj architektur
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
8 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Historick´y pˇrehled vybran´ych grafick´ych karet N´ azev GeForce GTX 480 GeForce GTX 680 Tesla M2090 Tesla K20 Quadro 2000D Quadro K5000
Rok 2010 2012 2011 2012 2010 2012
N´ azev GeForce GTX 480 GeForce GTX 680 Tesla M2090 Tesla K20 Quadro 2000D Quadro K5000
CC 2.0 3.0 2.0 3.5 2.1 3.0
Poˇcet jader 480 1536 512 2496 192 1536
Rychlost pamˇeti 177.4 GB/s 192.2 GB/s 177 GB/s 208 GB/s 41.6 GB/s 173 GB/s
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
Pamˇet’ 1.5GB GDDR5 2GB GDDR5 6 GB GDDR5 5 GB GDDR5 1 GB GDDR5 4GB GDDR5
V´ykon SP 1.345 TFlops 3.09 TFlops 1.33 TFlops 3.52 TFlops ? 2.1 TFlops 9 / 30
Rozhrann´ı 384-bit 256-bit 384-bit 384-bit 128-bit 256-bit
Cena 6 000 Kˇc 12 000 Kˇc 77 000 Kˇc 100 000 Kˇc 10 000 Kˇc 40 000 Kˇc ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Klasifikace GPU Flynnova taxonomie - Klasifikace poˇc´ıtaˇcov´ych architektur Single Data Multiple Data
Single Instruction SISD SIMD
Multiple Instruction MISD MIMD
ˇ Rady karet NVIDIA • GeForce - pomal´ a ve dvojit´e pˇresnosti • Tesla - profesion´ aln´ı, rychl´a dvojit´a pˇresnost • Quadro - urˇ cena sp´ıˇse pro grafiku
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
10 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Porovn´an´ı v´ykonu
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
11 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Architektura GPU
Fermi • SM - Multiprocesor s CUDA
j´adry • SFU - Zpracov´ an´ı sloˇzit´ych
operac´ı • LD/ST - Naˇ c´ıt´an´ı dat
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
12 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Typy pamˇet´ı
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
13 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Architektura vl´aken Device Global memory
Grid
Urˇcen´ı architektury // Urceni dimenzi dim3 gridD ( gX , gY , gZ ) ; dim3 blockD ( bX , bY , bZ ) ; // Volani kernelu kernel < < < gridD , blockD > > >( parametry ) ;
Block[0,0]
Block[0,1]
Block[0,2]
Block[0,3]
Shared memory
Shared memory
Shared memory
Shared memory
Block[1,0]
Block[1,1]
Block[1,2]
Block[1,3]
Shared memory
Shared memory
Shared memory
Shared memory
Block[1,2] Thread[0,0] R e g i s t e r
R e g i s t e r
Thread[1,0] R e g i s t e r
Thread[1,1] R e g i s t e r
Thread[2,0] R e g i s t e r
Thread[0,1]
Thread[2,1] R e g i s t e r
Thread[0,2] R e g i s t e r
Thread[1,2] R e g i s t e r
Thread[2,2] R e g i s t e r
Thread[0,3] R e g i s t e r
Thread[1,3] R e g i s t e r
Thread[2,3] R e g i s t e r
Thread[0,4] R e g i s t e r
Thread[1,4] R e g i s t e r
Thread[2,4] R e g i s t e r
Thread[0,5] R e g i s t e r
Thread[1,5] R e g i s t e r
Thread[2,5] R e g i s t e r
Shared memory
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
14 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Pˇriˇrazov´an´ı blok˚ u na SM
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
15 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Automatick´e promˇenn´e Dodefinovan´a struktura dim3 • blockDim.x - Obsahuje poˇ cet vl´aken uvnitˇr bloku • gridDim.x - Obsahuje poˇ cet blok˚ u uvnitˇr mˇr´ızky • blockIdx.x - Obsahuje index rezidentn´ıho bloku • threadIdx.x - Obsahuje index vl´ akno v bloku
Serializace index˚ u int idx = blockDim . x * blockIdx . x + threadIdx . x ;
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
16 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Jednoduch´y pˇr´ıklad C++ 1 2 3 4 5 6
1 2 3 4 5 6 7 8 9 10 11
__global__ void fillV ( float *v , float c ) { int idx = blockDim . x * blockIdx . x + threadIdx . x ; if ( idx < L ) { v [ idx ] = c ; } }
int main ( void ) { float hV [ L ]; float * dV ; for ( int i = 0; i < L ; i ++) hV [ i ] = 0.0; cudaMalloc (( void **) & dV , sizeof ( float ) * L ) ; cudaMemcpy ( dV , & hV , sizeof ( float ) *L , c u d a M e m c p y H o s t T o D e v i c e ) ; fillV < < <1 ,L > > >( dV , 99.1) ; cudaMemcpy (& hV , dV , sizeof ( float ) *L , c u d a M e m c p y D e v i c e T o H o s t ) ; for ( int i = 0; i < L ; i ++) printf ( " % f \ n " , hV [ i ]) ; return 0; }
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
17 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Optimalizace Pˇredch´azet rozdˇelen´ı (Divergenci) warpu • Vznik´ a napˇr. kv˚ uli podm´ınk´am typu: if ( threadIdx . x < 11) { branch1 () ; } else { branch2 () ; }
• Doch´ az´ı k serializaci • Potˇreba prov´ est obˇe vˇetve
ˇ ast vl´aken warpu nepracuje, spust´ı se v dalˇs´ı iteraci • C´
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
18 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Knihovny
Uˇziteˇcn´e knihovny • CUBLAS - CUDA implementace knihovny BLAS • CUSPARSE - CUDA implementace operac´ı s ˇr´ıdk´ ymi maticemi • MAGMA - CUDA implementace knihovny LAPACK • Thrust - Paraleln´ı algoritmy pro GPU, efektivnˇ ejˇs´ı
programov´an´ı
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
19 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
CUBLAS - porovn´an´ı v´ykonu
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
20 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Maticov´e n´asoben´ı na CPU (GLS BLAS) 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18
int main ( void ) { float * h_A , * h_B , * h_C ; int n2 = N * N ; float alfa = 1.0 , beta = 0.0; // * // I n i t i a l i z e host matrices h_A = new float [ n2 ]; h_B = new float [ n2 ]; h_C = new float [ n2 ]; // Compute cblas sgemm cblas_sgemm ( CblasColMajor , CblasNoTrans , CblasNoTrans , N , N , N , alfa , h_A , N , h_B , N , beta , h_C , N ) ; // ** return 0; }
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
21 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Maticov´e n´asoben´ı na GPU (CUBLAS) 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27
// * float * d_A = 0 , * d_B = 0 , * d_C = 0; cubl asStatus_t status ; cublasHandle_t handle ; // ** // I n i t i a l i z e device matrices C U D A _ C HEC K_R ETU RN ( cudaMalloc (( void **) & d_A , n2 * sizeof ( float ) ) ) ; C U D A _ C HEC K_R ETU RN ( cudaMalloc (( void **) & d_B , n2 * sizeof ( float ) ) ) ; C U D A _ C HEC K_R ETU RN ( cudaMalloc (( void **) & d_C , n2 * sizeof ( float ) ) ) ; // T r ansfer data into global memory and fill zeros into h_C c ub la sSetMatrix (N , N , sizeof ( float ) , h_A , N , d_A , N ) ; c ub la sSetMatrix (N , N , sizeof ( float ) , h_B , N , d_B , N ) ; C U D A _ C HEC K_R ETU RN ( cudaMemset ( d_C , 0 , n2 * sizeof ( float ) ) ) ; /* I n i t i a l i z e CUBLAS */ status = cublasCreate (& handle ) ; /* P e rforms o p e r a t i o n using cublas */ cublasSgemm ( handle , CUBLAS_OP_N , CUBLAS_OP_N , N , N , N , & alfa , d_A , N , d_B , N , & beta , d_C , N ) ; // Clear memory C U D A _ C HEC K_R ETU RN ( cudaFree ( d_A ) ) ; C U D A _ C HEC K_R ETU RN ( cudaFree ( d_B ) ) ; C U D A _ C HEC K_R ETU RN ( cudaFree ( d_C ) ) ; C U D A _ C HEC K_R ETU RN ( cudaDeviceReset () ) ; delete [] h_A ; delete [] h_B ; delete [] h_C ;
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
22 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Neuronov´a s´ıt’ s pˇrep´ınac´ımi jednotkami
´ cel Uˇ • Separaˇ cn´ı u ´lohy • Aproximace funkc´ı
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
23 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Stavebn´ı kameny NNSU Neuron s pˇrep´ınac´ı jednotkou (NSU) • Pˇrep´ınac´ı jednotka • Perceptrony • Vstupy rozdˇ elov´any na pereptrony pomoc´ı NSU
Blok • Zˇretˇ ezen´ı NSU • Bloky uspoˇr´ ad´any do acyklick´eho grafu
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
24 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Architektura NNSU - zn´azornˇen´ı NNSU
B4
NSU2 v
NSU1
Vstup
SU
v
v
B1
v
B2
NSU2
v
v N2
N1
v
v
v
v
B4
B3
NSU3
vv
v
v
v v Výstup
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
25 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Serializace s´ıtˇe
IP k´od • Reprezentace s´ıtˇ e - Program Symbol Tree’s • Serializace - Read˚ uv k´ od • Kombinace PST a Readov´ ych k´ od˚ u - IPCode
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
26 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Uˇcen´ı s´ıtˇe Prvn´ı u´roveˇn • Architektura stromov´ e struktury blok˚ u • Genetick´ a optimalizace • Paraleln´ı implementace
Druh´a u´roveˇn • Uˇ cen´ı perceptron˚ u kaˇzd´eho NSU
ˇ sen´ı soustav line´arn´ıch rovnic • Reˇ • S´ eriov´a implementace
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
27 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Integrace CUDA do projektu
´ Uprava Make soubor˚ u • Pˇredpis pro pˇreklad .cu soubor˚ u • Nastaven´ı cest ke knihovn´ am • Vytvoˇren´ı obalov´ eho k´ odu pro CUDA
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
28 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Dalˇs´ı kroky
Urˇcit pr´ah velikosti u´lohy • Mal´ eu ´lohy spouˇstˇet na CPU • Rozs´ ahlejˇs´ı u ´lohy spouˇstˇet na GPU
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
29 / 30
ˇ V. Spanihel
´ Uvod
Hardware
Knihovny
NNSU
Konec
Prostor pro dotazy
Vyuˇzit´ı GPU k urychlen´ı uˇ cen´ı neuronov´ ych s´ıt´ı
30 / 30
ˇ V. Spanihel
