Introduzione al Calcolo Parallelo
GPGPU – CUDA
Girolamo Giudice
Seminario di Bioinformatica
Introduzione al calcolo parallelo
•
•
•
•
•
•
Cenni sul calcolo sequenziale
Cenni sul calcolo parallelo
Perché usare il calcolo parallelo
Architettura hardware GPU - CUDA
Modello Software Cuda
Esempio pratico
Introduzione al calcolo parallello
• Benchmark di alcuni tool Bioinformatici
• Vento sulla GPU
Evoluzione della CPU
• Negli ultimi 20 anni i microprocessori basati su una
singola CPU hanno avuto un rapido incremento
nelle prestazioni e una diminuzione dei costi.
• Questa corsa ha subito una battuta d’ arresto a
causa dei consumi e dei problemi di riscaldamento
• 15 nov 2004 p4 3,8ghz
• 28 mag 2011 I7extreme 3,6ghz
Evoluzione della CPU
I produttori di microprocessori si sono orientati
verso modelli con più unità di processo (multi
core),allo scopo di aumentare la potenza di
calcolo.
Intel ha presentato un 80 core
Problemi dei Multi-core
• Tradizionalmente i programmi sono stati scritti
per essere eseguiti su un computer con una
singola CPU ( modello Von Neuman).
• La stragrande maggioranza delle applicazioni
sono costituite da programmi sequenziali
• I processori dual core sono praticamente lo
standard attuale
Cenni di calcolo sequenziale
• Un problema viene suddiviso in sequenze discrete di istruzioni
che vengono eseguite (di solito) una dopo l’altra
• In un dato istante di tempo solo una istruzione è in
esecuzione sulla CPU
Cenni di calcolo parallelo
• Il calcolo parallelo è l’uso di più unità di
computazione ( CPU multi core o multi CPU)
per risolvere problemi
• Storicamente è stato sempre un paradigma
costoso e di alto livello
Cenni di calcolo parallelo
• Il calcolo viene eseguito su più CPU o su CPU multicore
o dual thread
• Il problema viene decomposto in componenti discrete
che possono essere eseguite concorrentemente
• Le istruzioni sono eseguite simultaneamente su CPU
differenti
Tassonomia di Flynn
Tassonomia di Flynn
SISD
SIMD
MISD
MIMD
Perché usare il calcolo parallelo
• Risolvo un problema più grande nello stesso
tempo (SCALE – UP)
• Lo stesso problema in minor tempo (SPEEDUP)
• Contenere i costi
• Sfruttare meglio la RAM
• Aumentare l’affidabilità
• Utilizzare risorse distribuite
GPGPU / CUDA
• GPGPU: utilizzare il processore della scheda grafica
(GPU) per scopi diversi dalla tradizionale creazione di
un’immagine tridimensionale.
• Le GPU sono processori multicore ad elevate
prestazioni, il loro avvento è relativamente recente.
• Le prime soluzioni programmabili risalgono al
2006,precedentemente erano dedicate solo allo
sviluppo della grafica e dei videogiochi.
• Le GPU sono diventate processori paralleli general
purpose con interfacce di programmazione con
supporto ai linguaggi di programmazione come il C.
Differenze Macroscopiche CPU / GPU
Architettura CUDA G80
Host = CPU
Device = GPU
1
2
3
4
C: Compute
U: Unified
D: Device
A: Architecture
5
6
7
8
9
10
11 12
13 14
15 16
Architettura Cuda
Ciascun Streaming Multiprocessor contiene al suo interno:
• 8 Stream Processor (add sub,mul su int e float)
• SFU(super funciton unit): seno,coseno,log,inv,exp
• Shared memory per tutti i thread in esecuzione sul SM
• Cache per dati e istruzioni
• Unità per la decodifica delle istruzioni (decodifica una istruzione
ogni 4 cicli di clock)
Architettura hardware
• Mascherare la latenza della memoria globale con migliaia di thread
• Struttura di memoria semplice ma a bassa latenza, anziché ad
accesso ottimizzato , ma complesso
• Nessuna priorità sui thread
• No context switch
• No overhead
• SIMT (single instruction multiple thread) tutti i thread eseguono la
stesso istruzione ma su dati diversi
Cuda: Modello di esecuzione
Un codice Cuda alterna porzioni di codice
seriale, eseguito dalla CPU e di codice
parallelo eseguito dalla GPU.
Le porzioni di codice eseguite sulla GPU
sono note come kernel (~ funzione in
C/C++)
Il kernel, è definito come una griglia di
blocchi che vengono assegnati ai vari
multiprocessori, e rappresentano un
parallelismo a grana grossa.
Ogni
blocco
esegue
l’unità
di
computazione fondamentale, il thread.
Un thread può appartenere ad un solo
blocco ed è univocamente identificato da
un ID.
Multidimensionalità degli IDs
Il codice parallelo viene lanciato,
dalla CPU, sulla GPU , questa esegue
un solo kernel alla volta.
La dimensione della griglia si misura
in blocchi questi possono essere:
Block: 1-D o 2-D
La dimensione dei blocchi si misura
in thread
Thread 1-D,2-D,3-D
Cuda memory model
Tipi di memoria
• Global (device) memory
(R/W)
• Shared memory (R/W)
• Registers (R/W locale per
thread)
• Constant (R/O)
• Texture (R/O)
Global,costant e texture memory
sono persistenti a diversi lanci di
kernel
Si minimizza il transfer rate
bottleneck
Classi di applicazioni
• Presenza di molte operazioni
matematiche(grande intensità aritmetica)
• Elevato grado di parallelismo (le stesse
operazioni vengono ripetute per una grande
quantità di dati)
• Condizioni di controllo limitate
• Minima dipendenza tra i dati
Linguaggi che supportano cuda
Esempio
• 2 vettori da 100.000 elementi
• Su ogni elemento del vettore dobbiamo eseguire
questa operazione log(h_a[i]*h_b[i])
• Quanti blocchi?
• Quanti thread?
Fissiamo per esempio 512 thread
Dimensione del blocco = 100000/512=195.31
Arrotondiamo a 196
n°thread=196*512=100352
#include <stdio.h>
// implementazione del kernel
__global__ void Kernel(float *d_a,float *d_b,float *d_c)
{
// calcolo dell'indice di thread
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if(idx<100000)
d_c[idx] =log(d_a[idx]*d_b[idx]);
}
// Dichiariamo il main
int main( int argc, char** argv)
{
int n=100000;
time_t begin,end;
// puntatore per la struttura dati sull'host
float *h_a=(float*) malloc(sizeof(float)*n);
float *h_b=(float*) malloc(sizeof(float)*n);
float *h_c=(float*) malloc(sizeof(float)*n);
//inizializzo il vettore numeri casuali
for(int i=0;i<n;i++)
{
h_a[i]=rand();
h_b[i]=rand();
}
begin = clock();
for(int i=0;i<n;i++)
h_c[i] =log(h_a[i]*h_b[i]);
end=clock();
float time_cpu = (double)(end-begin)/CLOCKS_PER_SEC;
printf("CPU time %.20lf\n",time_cpu);
// puntatore per la struttura dati sul device
float *d_a=NULL;
float *d_b=NULL;
float *d_c=NULL;
//verifico al secondo lancio del kernel
for(int i=0;i<2;i++)
{
begin = clock();
//malloc e memcopy host to device
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
cudaMemcpy( d_a, h_a, sizeof(float)*n,
cudaMemcpyHostToDevice) ;
cudaMemcpy( d_b, h_b, sizeof(float)*n,
cudaMemcpyHostToDevice) ;
cudaMemcpy( d_c, h_c, sizeof(float)*n,
cudaMemcpyHostToDevice) ;
// definizione della grandezza della griglia e dei blocchi
int numBlocks = 196;
int numThreadsPerBlock = 512;
// Lancio del kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
Kernel<<< dimGrid, dimBlock >>>( d_a,d_b,d_c );
// blocca la CPU fino al completamento del kernel sul device
cudaThreadSynchronize();
// Esegue la copia dei risultati dalla memoria del device a quella dell'host
cudaMemcpy( h_c, d_c, n, cudaMemcpyDeviceToHost );
end = clock();
}
float time_gpu = (double)(end-begin)/CLOCKS_PER_SEC;
printf("GPU time %.20lf\n",time_gpu);
// libera la memoria sul device
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
// libera la memoria sull'host
free(h_a);
free(h_b);
free(h_c);
return 0;
}
Inizialmente:
float *h_a=(float*) malloc(sizeof(float)*n);
float *h_b=(float*) malloc(sizeof(float)*n);
float *h_c=(float*) malloc(sizeof(float)*n);
GPU
CPU
Array h_a
Host’s
memory
Array h_b
Array h_c
Device’s
memory
Allocare memoria sulla GPU
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
cudaMalloc( (void**) &d_a, sizeof(float)*n) ;
GPU
CPU
Host’s
memory
Array h_a
Array d_a
Array h_b
Array d_b
Array h_c
Array d_c
Device’s
memory
Copiare il contenuto dalla host memory
alla device memory
cudaMemcpy( d_a, h_a, sizeof(float)*n, cudaMemcpyHostToDevice) ;
cudaMemcpy( d_b, h_b, sizeof(float)*n, cudaMemcpyHostToDevice) ;
cudaMemcpy( d_c, h_c, sizeof(float)*n, cudaMemcpyHostToDevice) ;
CPU
Host’s
memory
GPU
Array h_a
Array d_a
Array h_b
Array d_b
Array h_c
Array d_c
Device’s
memory
Eseguire il contenuto sulla GPU
__global__ void Kernel( float *d_a,float *d_b,float *d_c)
{
// calcolo dell'indice di thread
Kernel<<< 196, 512 >>>( d_a,d_b,d_c );
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if(idx<100000)
d_c[idx] =log(d_a[idx]*d_b[idx]);
}
GPU MPs
GPU
CPU
Host’s
memory
Array h_a
Array d_a
Array h_b
Array d_b
Array h_c
Array d_c
Device’s
memory
In the GPU
Thread
0
Thread
1
Thread
2
…
Thread
512
Thread
0
Thread
1
Thread
2
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
…
…
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
Block 0
d_c[idx]=
Log(d_a[idx]
+d_b[idx]
… …
……
Block 196
Thread
512
Restituire il risultato
cudaMemcpy( h_c, d_c, n, cudaMemcpyDeviceToHost );
Array h_c
Array d_c
Host’s Memory
Tempi
CPU 0.01
GPU 0.002
100.000 el
GPU Card’s Memory
Tempi
CPU 0.1
GPU 0.007
1.000.000 el
Tempi
CPU 0.8
GPU 0.037
10.000.000 el
Applicazioni tipiche
•
•
•
•
•
•
•
•
•
•
•
Elaborazione video
Astrofisica
Finanza
Fisica di gioco
Modellazione fisica
Analisi numerica
DSP
Imaging medicale
Data mining
Dinamica molecolare
Bioinformatica
http://gpu.epfl.ch/sw.html
Mcode:Finds clusters in a network
Subgraph isomorphism
Subgraph matching.
Un grafo G(V,E) e un sottografo isomorfo a G1(V1,E1) se esiste una funzione iniettiva f:VV1 tale che (u,v)E se e solo se
(f(u),f(v))E1.
La ricerca di sottostrutture all’interno di un grafo target è un processo estremamente
oneroso dal punto di vista computazionale (problema NP-completo).
Il processo di ricerca di una query si articola in tre fasi
• Preprocessing
• Filtering
• Matching
Esempio
#Graph
4
71
83
27
44
3
01
02
23
71
0
27
2
1
83
44
3
Esempio
#graph
15
69
72
37
22
97
95
50
8
88
8
50
69
12
81
36
10
16
07
01
71
10 1
11 10
11 13
11 6
11 3
11 9
11 8
11 5
11 2
11 4
11 12
14 11
69
9
7
8
88
8
50
4
97
1
72
8
11
0
95
69
3
37
5
22
2
13
6
36
50
14
81
12
12
Grafo Query
#Query
3
69
50
8
2
01
02
69
0
8
2
1
50
Preprocessing
Nodo iniziale / Nodo finale
#graph
16
8
69
8
72
69
72
50
72
69
50
69
81
69
50
69
22
69
8
69
88
69
95
69
37
69
97
69
12
36
69
Nodo iniziale / Nodo finale
#Query
2
69
50
69
8
Applichiamo CUDA
Nodo iniziale / Nodo finale
#graph
16
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
8
8
69
50
69
69
69
69
69
69
69
69
69
69
36
69
72
72
72
50
81
50
22
8
88
95
37
97
12
69
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
Nodo iniziale / Nodo finale
#Query
2
69
50
69
8
Nodo iniziale / Nodo finale
#graph
16
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
8
8
69
50
69
69
69
69
69
69
69
69
69
69
36
69
72
72
72
50
81
50
22
8
88
95
37
97
12
69
1° kernel
69
50
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
Nodo iniziale / Nodo finale
#graph
16
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
69
8
69
50
69
69
69
69
69
69
69
69
69
69
36
8
72
72
72
50
81
50
22
8
88
95
37
97
12
69
2° kernel
69
8
Tid 1
Tid 2
Tid 3
Tid 4
Tid 5
Tid 6
Tid 7
Tid 8
Tid 9
Tid 10
Tid 11
Tid 12
Tid 13
Tid 14
Tid 15
Tid 16
Pruning
Foresta di grafi
50
Query
69
0
1
8
8
50
69
8
2
50
69
Ricapitolando
CUDA
2° pruning
VF2
Risultati
Rete utilizzata:
Scalefree2000
Composta da 2000 nodi e 3997 archi
Query Test:
Query4
Query16
Query64
Hardware utilizzato:
Intel Core 2 duo E4400 (2 GHz)
Nvidia Geforce Gts 250 (128 Cuda cores)
Vento cpu VS gpu query4
0.035
0.03
Tempo(s)
0.025
0.02
0.015
0.01
0.005
0
ventoCPU
query query query query query query query query query query query query query query query query query query query query query
4_00 4_01 4_01 4_01 4_02 4_02 4_03 4_03 4_04 4_04 4_04 4_05 4_05 4_06 4_06 4_07 4_07 4_08 4_08 4_09 4_09
5
1
5
8
3
7
2
6
1
4
8
3
8
3
7
2
8
2
7
3
7
0
0
0.02
0
0
0
0.03 0.02 0.02 0.01
0
0
0
0
0
0
0
0
0
0
0
ventoGPU 0.01 0.01 0.001 0.005 0.006 0.004 0.005 0.001 0.01 0.004 0.005 0.006 0.006 0.027 0.004 0.006 0.005 0.005 0.006 0.006 0.004
Vento CPU VS GPU query16
0.1
0.09
0.08
Tempo(s)
0.07
0.06
0.05
0.04
0.03
0.02
0.01
0
query query query query query query query query query query query query query query query query query query query query
16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0 16_0
01
04
08
12
16
20
24
29
33
44
49
58
67
71
75
79
85
90
94
98
ventoCPU 0.01
0
0.01
0
0.02
0
0
0
0
0
0
ventoGPU 0.09 0.001 0.016 0.015 0.005 0.008 0.006 0.008 0.008 0.006 0.001
0
0
0
0
0.02 0.001
0
0
0
0.02 0.01
0
0
0.001
0
0.001 0.006 0.008
Vento cpu/gpu query64
12
10
8
6
4
2
0
query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6 query6
4_004 4_009 4_014 4_024 4_027 4_031 4_036 4_041 4_049 4_055 4_059 4_062 4_065 4_078 4_082 4_087 4_091 4_095
vento cpu 0.08
1.12
0.02
0.44
1.79
2.95
0.23
0.28
9.7
0.03
0.05
0.39
6.24
0.02
ventogpu
0.166
0.04
0.08
0.22
0.038
0.05
0.068
0.96
0.018
0.02
0.03
0.32
0.02 0.0199 0.024 0.1426 0.82
0.12
0.02
0.01
0.55
7.55
Vento CPU/GPU query64
300
250
Tempi(s)
200
150
100
50
0
query64_019
query64_074
query64_069
ventocpu
281.52
192.29
91.43
ventogpu
28.22
39.83
15.44
graph reduced
0.021
0.019
0.016
Rete utilizzata:
YeastNetworkRand
Composta da 5589 nodi e 92835 archi
Query Test:
Query4
Query8
Query16
Query32
Query64
Query128
Hardware utilizzato:
Intel Core 2 duo E4400 (2 GHz)
Nvidia Geforce Gts 250 (128 Cuda cores)
VentoCPU vs VentoGPU
0.3
0.25
Tempi(s)
0.2
0.15
0.1
0.05
0
query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_ query4_
001
010
019
032
037
040
043
046
050
052
054
060
067
071
079
082
090
ventoCPU
0.03
0.02
0.02
0.14
0
0.02
0.02
0.02
0.02
0.02
0.02
0.02
0.02
0.05
0.28
0.02
0.03
ventoGPU
0.01
0.019
0.007
0.017
0.007
0.012
0.008
0.007
0.007
0.007
0.027
0.018
0.007
0.026
0.061
0.027
0.018
Query 8
70
60
Circa 60x più rapido
50
40
ventoCPU
ventoGPU
30
20
Circa 6x più rapido
10
0
query8_099
query8_018
VentoCPU vs VentoGPU
0.25
0.2
Tempi(s)
0.15
0.1
0.05
0
query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16 query16
_005
_012
_021
_026
_029
_035
_053
_056
_061
_068
_071
_074
_079
_083
_096
_098
ventoCPU
0.01
0.23
0.09
0.01
0.17
0.14
0.05
0.01
0.12
0.09
0.01
0.02
0.23
0.19
0.02
0.11
ventoGPU
0.013
0.0313
0.0213
0.015
0.021
0.0213
0.021
0.011
0.01
0.01
0.012
0.01
0.01
0.011
0.021
0.012
Query 16
Guadagno da 6x a 164x
Query 16
60
5136x più rapido
50
40
ventoCPU
30
ventoGPU
20
10
Circa 26x più rapido
20x più rapido
0
query16_002
query16_087
query16_090
Query 32
VentoCPU vs VentoGPU
0.5
0.45
0.4
0.35
Tempi(s)
0.3
0.25
0.2
0.15
0.1
0.05
0
query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3 query3
2_003 2_009 2_021 2_022 2_027 2_036 2_045 2_046 2_053 2_062 2_068 2_072 2_074 2_085 2_095 2_098 2_099
VentoCPU
0.25
0.08
0.39
ventoGPU
0.13
0.01
0.025 0.014 0.016 0.026 0.016 0.026 0.035 0.016 0.0001 0.015 0.016 0.016 0.015 0.016 0.015
0.03
0.06
0.05
0.03
0.05
0.12
0.05
0.14
0.11
0.44
0.17
0.22
0.28
0.06
3.5
199x più rapido
3
106x più rapido
2.5
2
VentoCPU
VentoGPU
1.5
24x più rapido
30x più rapido
6x più rapido
1
0.5
0
query32_026
query32_013
query32_015
query32_061
query32_067
80
28x più rapido
70
50
40
30
20
10
VentoCPU impiega oltre 7 ore
60
ventoCPU
VentoGPU
24x più rapido
0
query32_035
query32_051
query32_091
VentoCPU vs VentoGPU
1.6
1.4
1.2
Tempi(s)
1
0.8
0.6
0.4
0.2
0
query128 query128 query128 query128 query128 query128 query128 query128 query128 query128 query128 query128 query128 query128
_006
_022
_029
_044
_056
_059
_065
_075
_081
_084
_088
_094
_096
_098
VentoCPU
0.22
0.12
0.09
0.19
0.17
0.42
0.03
0.14
0.55
0.08
0
0.14
0.44
1.51
VentoGPU
0.09
0.04
0.08
0.053
0.049
0.07
0.037
0.06
0.1
0.037
0.047
0.1
0.047
0.055
35
10,1x più rapido
30
25
20
VentoCPU
172x più rapido
VentoGPU
15
10
5
26x più rapido
9,26x più rapido
16,2x più rapido
16,6x più rapido
0
query128_001
query128_041
query128_049
query128_052
query128_063
query128_078
• L’elaborazione della query128_036,da parte di
ventoCPU, è stata interrotta dopo oltre 7 ore
di elaborazione
Confronto con SING
• Rete:ScaleFree_2000_128
• n° Query4 : 100
• n° Match Effettivi: 31
• Query da valutare con VentoGPU:32
• Falsi positivi VentoGPU:1
• Query da valutare con Sing: 99
• Falsi positivi Sing: 68
• Rete:ScaleFree_2000_128
• n° Query16 : 100
• n° Match Effettivi: 49
• Query da valutare con VentoGPU:49
• Falsi positivi VentoGPU:0
• Query da valutare con Sing: 50
• Falsi positivi Sing: 1
• Rete:ScaleFree_2000_128
• n° Query64 : 100
• n° Match Effettivi: 40
• Query da valutare con VentoGPU:40
• Falsi positivi VentoGPU:0
• Query da valutare con Sing: 42
• Falsi positivi Sing: 2
• Rete:Yeastnetworkrand
• n° Query4 : 100
• n° Match Effettivi: 68
• Query da valutare con VentoGPU:75
• Falsi positivi VentoGPU:7
• Rete:Yeastnetworkrand
• n° Query16 : 100
• n° Match Effettivi: 35
• Query da valutare con VentoGPU:61
• Falsi positivi VentoGPU:26
• Rete:Yeastnetworkrand
• n° Query64 : 100
• n° Match Effettivi: 34
• Query da valutare con VentoGPU: 60
• Falsi positivi VentoGPU:26
Questions?