Slides

Le schede grafiche

CUDA

Implementazione

Risultati

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture) Stefano Brilli Universit` a degli studi di Firenze

Ottobre 2008

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

CUDA

CUDA `e una tecnologia di NVIDIA Corporation che permette di sfruttare l’architettura parallela delle moderne schede grafiche, per sviluppare algoritmi di calcolo altamente efficienti.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Argomenti trattati

1

Le schede grafiche

2

CUDA

3

Implementazione

4

Risultati

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Argomenti trattati

1

Le schede grafiche

2

CUDA

3

Implementazione

4

Risultati

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Eseguono il rendering di scene tridimensionali, per applicazioni e giochi per computers. Sono dotate di una propria memoria ed una propria unit`a di elaborazione: la Graphics Processing Unit (GPU). Le GPU si sono evolute fino a diventare potenti processori paralleli multicore

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Eseguono il rendering di scene tridimensionali, per applicazioni e giochi per computers. Sono dotate di una propria memoria ed una propria unit`a di elaborazione: la Graphics Processing Unit (GPU). Le GPU si sono evolute fino a diventare potenti processori paralleli multicore

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Eseguono il rendering di scene tridimensionali, per applicazioni e giochi per computers. Sono dotate di una propria memoria ed una propria unit`a di elaborazione: la Graphics Processing Unit (GPU). Le GPU si sono evolute fino a diventare potenti processori paralleli multicore

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

La maggior parte della loro superficie `e dedicata ad unit`a di calcolo, pittusto che al data caching e al controllo di flusso.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Argomenti trattati

1

Le schede grafiche

2

CUDA

3

Implementazione

4

Risultati

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

CUDA

CUDA `e un ambiente di sviluppo che estende il linguaggio C. Permette di definire dei kernel. Ovvero particolari funzioni che verranno eseguite dalla GPU. Offre un modello di programmazione basato su: 1 2 3

Stefano Brilli

Gerarchia dei thread Memoria condivisa Sincronizzazione a barriera

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

CUDA

CUDA `e un ambiente di sviluppo che estende il linguaggio C. Permette di definire dei kernel. Ovvero particolari funzioni che verranno eseguite dalla GPU. Offre un modello di programmazione basato su: 1 2 3

Stefano Brilli

Gerarchia dei thread Memoria condivisa Sincronizzazione a barriera

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

CUDA

CUDA `e un ambiente di sviluppo che estende il linguaggio C. Permette di definire dei kernel. Ovvero particolari funzioni che verranno eseguite dalla GPU. Offre un modello di programmazione basato su: 1 2 3

Stefano Brilli

Gerarchia dei thread Memoria condivisa Sincronizzazione a barriera

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Gerarchia dei Thread

Multithreading Ogni kernel viene eseguito da migliaia di thread contemporaneamente. Gerarchia I Thread sono raggruppati in blocchi. L’insieme di tutti i blocchi forma la griglia di esecuzione.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Memoria condivisa I thread hanno accesso a diversi spazi di memoria Lettura/Scrittura Memoria Locale Memoria Condivisa Memoria Globale

Sola Lettura Memoria Costante Memoria delle Texture

. Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Memoria condivisa I thread hanno accesso a diversi spazi di memoria Lettura/Scrittura Memoria Locale Memoria Condivisa Memoria Globale

Sola Lettura Memoria Costante Memoria delle Texture

. Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Memoria condivisa I thread hanno accesso a diversi spazi di memoria Lettura/Scrittura Memoria Locale Memoria Condivisa Memoria Globale

Sola Lettura Memoria Costante Memoria delle Texture

. Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Sincronizzazione a barriera CUDA fornisce una primitiva di sincronizzazione a barriera: syncthreads() Ferma il thread che l’ha invocata, finch`e tutti gli altri thread dello stesso blocco non l’hanno invocata a loro volta. Problema Non esiste una primitiva di sincronizzazione globale! Soluzione Si attende il ritorno dalla chiamata ad un kernel: A questo punto tutti i thread hanno svolto il proprio compito.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Sincronizzazione a barriera CUDA fornisce una primitiva di sincronizzazione a barriera: syncthreads() Ferma il thread che l’ha invocata, finch`e tutti gli altri thread dello stesso blocco non l’hanno invocata a loro volta. Problema Non esiste una primitiva di sincronizzazione globale! Soluzione Si attende il ritorno dalla chiamata ad un kernel: A questo punto tutti i thread hanno svolto il proprio compito.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Sincronizzazione a barriera CUDA fornisce una primitiva di sincronizzazione a barriera: syncthreads() Ferma il thread che l’ha invocata, finch`e tutti gli altri thread dello stesso blocco non l’hanno invocata a loro volta. Problema Non esiste una primitiva di sincronizzazione globale! Soluzione Si attende il ritorno dalla chiamata ad un kernel: A questo punto tutti i thread hanno svolto il proprio compito.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in C Ambiente C void vecAdd(float* A, float* B, float* C){ int i; for(i=0;i
A:

Stefano Brilli

a1 ↑

a2 ↑

a3 ↑

a4 ↑

a5 ↑

a6 ↑

...

alen ↑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in CUDA Ambiente CUDA __global__ void vecAdd(float* A, float* B, float* C){ int i = blockDim.x*blockIdx.x+threadIdx.x; for(;i>>(A, B, C); }

A:

Stefano Brilli

a1 ⇑

a2 ⇑

... ···

ak ⇑

ak+1 ⇑

ak+2 ⇑

. . . a2k ··· ⇑

... ⇑

alen ⇑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in CUDA Ambiente CUDA __global__ void vecAdd(float* A, float* B, float* C){ int i = blockDim.x*blockIdx.x+threadIdx.x; for(;i>>(A, B, C); } k = numero di thread totali. A:

Stefano Brilli

a1 ⇑

a2 ⇑

... ···

ak ⇑

ak+1 ⇑

ak+2 ⇑

. . . a2k ··· ⇑

... ⇑

alen ⇑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in CUDA Ambiente CUDA __global__ void vecAdd(float* A, float* B, float* C){ int i = blockDim.x*blockIdx.x+threadIdx.x; for(;i>>(A, B, C); } k = numero di thread totali. A:

Stefano Brilli

a1 ⇑

a2 ⇑

... ···

ak ⇑

ak+1 ⇑

ak+2 ⇑

... ···

a2k ⇑

... ⇑

alen ⇑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Somma membro a membro di due vettori in CUDA Ambiente CUDA __global__ void vecAdd(float* A, float* B, float* C){ int i = blockDim.x*blockIdx.x+threadIdx.x; for(;i>>(A, B, C); } k = numero di thread totali. A:

Stefano Brilli

a1 ⇑

a2 ⇑

... ···

ak ⇑

ak+1 ⇑

ak+2 ⇑

. . . a2k ··· ⇑

... ⇑

alen ⇑

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Argomenti trattati

1

Le schede grafiche

2

CUDA

3

Implementazione

4

Risultati

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Il Metodo dei Gradienti Coniugati Cos’`e Il metodo dei gradienti coniugati `e un metodo iterativo per la risoluzione di sistemi lineari. L’algoritmo d0 = b − Qx0 αk =

gkT dk dT k Qdk

xk+1 = xk + αk dk βk =

T gk+1 Qdk T dk Qdk

dk+1 = −gk+1 + βk dk Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Le implementazioni ccgmamma.c Implementazione C

−→ cgmamma.m Implementazione Matlab

−→

nvcgmamma.cu Implementazione CUDA

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Fasi dello sviluppo

1 2

Implementazione del prodotto matrice-vettore. Tecniche di ottimizzazione. Allineamento conveniente dei dati. Uso della memoria condivisa per lo scambio dei risultati intermedi.

3

Stefano Brilli

Implementazione delle restanti operazioni.

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Fasi dello sviluppo

1 2

Implementazione del prodotto matrice-vettore. Tecniche di ottimizzazione. Allineamento conveniente dei dati. Uso della memoria condivisa per lo scambio dei risultati intermedi.

3

Stefano Brilli

Implementazione delle restanti operazioni.

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Fasi dello sviluppo

1 2

Implementazione del prodotto matrice-vettore. Tecniche di ottimizzazione. Allineamento conveniente dei dati. Uso della memoria condivisa per lo scambio dei risultati intermedi.

3

Stefano Brilli

Implementazione delle restanti operazioni.

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Fasi dello sviluppo

1 2

Implementazione del prodotto matrice-vettore. Tecniche di ottimizzazione. Allineamento conveniente dei dati. Uso della memoria condivisa per lo scambio dei risultati intermedi.

3

Stefano Brilli

Implementazione delle restanti operazioni.

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Fasi dello sviluppo

1 2

Implementazione del prodotto matrice-vettore. Tecniche di ottimizzazione. Allineamento conveniente dei dati. Uso della memoria condivisa per lo scambio dei risultati intermedi.

3

Stefano Brilli

Implementazione delle restanti operazioni.

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Argomenti trattati

1

Le schede grafiche

2

CUDA

3

Implementazione

4

Risultati

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Performance

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Performance

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Rapporto prestazioni/prezzo

CPU Intel Core 2 Duo 2,66 GHz Costo: 150 e

GPU NVIDIA GeForce 8800 GTS 512 MHz

Costo:110 e

Risultati Prestazioni 20 volte superiori e costi inferiori rendono la GPU di NVIDIA 27 volte pi` u conveniente della CPU di Intel.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)

Le schede grafiche

CUDA

Implementazione

Risultati

Rapporto prestazioni/prezzo

CPU Intel Core 2 Duo 2,66 GHz Costo: 150 e

GPU NVIDIA GeForce 8800 GTS 512 MHz

Costo:110 e

Risultati Prestazioni 20 volte superiori e costi inferiori rendono la GPU di NVIDIA 27 volte pi` u conveniente della CPU di Intel.

Stefano Brilli

Universit` a degli studi di Firenze

Implementazione efficiente del Metodo dei Gradienti Coniugati in ambiente CUDA (Compute Unified Device Architecture)