Programmazione di processori massivamente paralleli: un approccio pratico

📚 Riepilogo del contenuto

Questo corso offre un'introduzione completa al calcolo GPU e alla programmazione parallela tramite l'ambiente CUDA C. Copre le architetture GPU, il parallelismo dati, la gestione dei thread, l'ottimizzazione della memoria e considerazioni avanzate sulle prestazioni, illustrate attraverso casi reali come la ricostruzione di immagini MRI e la visualizzazione molecolare.

Padroneggia l'arte del calcolo parallelo ad alte prestazioni con una guida pratica e basata su esercitazioni per CUDA e le architetture GPU.

Autore: David B. Kirk, Wen-mei W. Hwu

Ringraziamenti: Ian Buck, John Nickolls, team NVIDIA DevTech, Jensen Huang, David Luebke, Bill Bean, Simon Green, Mark Harris, Manju Hedge, Nadeem Mohammad, Brent Oster, Peter Shirley, Eric Young e Cyril Zeller.

🎯 Obiettivi di apprendimento

1. Distinguere tra filosofie progettuali e traiettorie di prestazioni di CPU multicore e GPU many-core.
2. Identificare i componenti chiave di un'architettura GPU moderna, inclusi i Streaming Multiprocessors (SM) e le strutture di memoria.
3. Applicare la legge di Amdahl per calcolare il guadagno teorico in velocità e identificare l'impatto dei colli di bottiglia sequenziali.
4. Confrontare le differenze architetturali tra pipeline fisse e array di processori programmabili unificate.
5. Spiegare il ruolo del "GPGPU" come passaggio intermedio e le limitazioni dei primi modelli di programmazione shader.
6. Analizzare come funzionalità hardware come operazioni atomiche, sincronizzazione barriera e supporto a precisione doppia hanno permesso il passaggio al calcolo generico scalabile.
7. Identificare e sfruttare il parallelismo dati all'interno degli algoritmi di moltiplicazione matrice-matrice.
8. Implementare la gestione della memoria del dispositivo, inclusa l'allocazione, il trasferimento dei dati tra host e dispositivo e la deallocazione.
9. Creare e lanciare kernel CUDA utilizzando indici thread appropriati e configurazioni di griglia/blocco.
10. Progettare gerarchie di thread multidimensionali (griglie e blocchi) per mappare strutture dati complesse sull'hardware GPU.

🔹 Lezione 1: Introduzione al calcolo parallelo e alle architetture GPU

Panoramica: Questa lezione esplora il cambiamento fondamentale dal calcolo sequenziale a quello parallelo, guidato dalle diverse filosofie progettuali di CPU e GPU. Gli studenti esamineranno le traiettorie "Multicore" vs "Many-core", comprendere la struttura hardware che consente alle GPU di raggiungere un throughput massiccio e imparare i vincoli matematici del guadagno di velocità tramite la legge di Amdahl.

Risultati dell'apprendimento:

• Distinguere tra filosofie progettuali e traiettorie di prestazioni di CPU multicore e GPU many-core.
• Identificare i componenti chiave di un'architettura GPU moderna, inclusi i Streaming Multiprocessors (SM) e le strutture di memoria.
• Applicare la legge di Amdahl per calcolare il guadagno teorico in velocità e identificare l'impatto dei colli di bottiglia sequenziali.

🔹 Lezione 2: Evoluzione e futuro del calcolo GPU

Panoramica: Questa lezione ripercorre il cammino architetturale dell'Unità di Elaborazione Grafica (GPU) dalla sua origine come hardware specializzato a funzione fissa per il rendering di triangoli fino allo stato attuale di potente processore parallelo unificato e general-purpose. Gli studenti esploreranno il passaggio da pipeline grafiche rigide a shader programmabili, l'emergere del movimento GPGPU e le architetture scalabili moderne che alimentano simulazioni scientifiche e ingegneristiche attuali.

Risultati dell'apprendimento:

• Confrontare le differenze architetturali tra pipeline fisse e array di processori unificati programmabili.
• Spiegare il ruolo del "GPGPU" come passaggio intermedio e le limitazioni dei primi modelli di programmazione shader.
• Analizzare come funzionalità hardware come operazioni atomiche, sincronizzazione barriera e supporto a precisione doppia hanno favorito il passaggio al calcolo generico scalabile.

🔹 Lezione 3: Struttura dei programmi CUDA e gestione della memoria

Panoramica: Questa lezione copre l'architettura fondamentale di un programma CUDA, sottolineando la distinzione tra esecuzione Host (CPU) e Device (GPU). Gli studenti impareranno a identificare il parallelismo dati nelle operazioni matriciali, gestire spazi di memoria separati tramite l'API CUDA e organizzare l'esecuzione parallela attraverso una gerarchia di griglie, blocchi e thread usando lo stile Single-Program, Multiple-Data (SPMD).

Risultati dell'apprendimento:

• Identificare e sfruttare il parallelismo dati negli algoritmi di moltiplicazione matrice-matrice.
• Implementare la gestione della memoria del dispositivo, inclusa l'allocazione, il trasferimento dei dati tra host e dispositivo e la deallocazione.
• Creare e lanciare kernel CUDA utilizzando indici thread e configurazioni di griglia/blocco appropriate.

🔹 Lezione 4: Thread e scheduling avanzati CUDA

Panoramica: Questa lezione esplora l'organizzazione gerarchica dei thread in CUDA, concentrandosi su come gli indici multidimensionali si mappino su dati fisici e risorse hardware. Dettaglia i meccanismi di sincronizzazione barriera e scalabilità trasparente, concludendo con i principi architetturali di assegnazione dei thread e schedulazione basata su warp utilizzati per ottenere tolleranza al ritardo nel calcolo ad alte prestazioni.

Risultati dell'apprendimento:

• Progettare gerarchie di thread multidimensionali (griglie e blocchi) per mappare strutture dati complesse sull'hardware GPU.
• Implementare un'indicizzazione precisa dei dati usando variabili CUDA incorporate (blockIdx, threadIdx, blockDim).
• Applicare la sincronizzazione barriera per garantire l'integrità dei dati mantenendo una scalabilità trasparente su diverse architetture GPU.

🔹 Lezione 5: Ottimizzazione della memoria e tiling della memoria condivisa

Panoramica: Questa lezione esplora come larghezza di banda della memoria e vincoli di risorse agiscano come colli di bottiglia principali nel calcolo parallelo. Dettaglia l'uso del "tiling" per ridurre il traffico della memoria globale e spiega il ruolo cruciale delle barriere di sincronizzazione (__syncthreads()) e della scelta strategica tra registri e memoria condivisa per ottimizzare le prestazioni.

Risultati dell'apprendimento:

• Analizzare come i limiti di registri e memoria condivisa determinino il livello di parallelismo (occupazione) in un kernel.
• Quantificare la riduzione del consumo di larghezza di banda della memoria globale ottenuta tramite tecniche di tiling.
• Identificare la necessità di funzioni di sincronizzazione per mantenere l'integrità dei dati durante l'accesso alla memoria condivisa.

🔹 Lezione 6: Analisi delle prestazioni e esecuzione SIMT

Panoramica: Questa lezione esplora i considerazioni architetturali e algoritmiche essenziali per ottimizzare i kernel CUDA. Passa dai modelli di esecuzione di base — specificamente l'unità Single-Instruction, Multiple-Thread (SIMT) e il partizionamento di warp — alle tecniche avanzate di ottimizzazione delle prestazioni, incluse il coalescing della memoria, la moltiplicazione matriciale tiling e il partizionamento dinamico delle risorse di Streaming Multiprocessor (SM).

Risultati dell'apprendimento:

• Analizzare la mappatura di blocchi di thread multidimensionali sull'ordine lineare di esecuzione dei warp hardware.
• Valutare e minimizzare la divergenza del flusso di controllo negli algoritmi di riduzione parallela.
• Ottimizzare la larghezza di banda della memoria globale implementando pattern di accesso coalesciato e tiling.

🔹 Lezione 7: Aritmetica floating-point e accuratezza numerica

Panoramica: Questa lezione copre l'architettura fondamentale dei numeri floating-point, con particolare attenzione ai componenti dello standard IEEE 754: segno, esponente codificato in eccesso e mantissa normalizzata. Gli studenti esploreranno come questi schemi binari si mappino su una linea numerica discreta e come i limiti di questa rappresentazione influenzino l'accuratezza di algoritmi complessi come somme su larga scala.

Risultati dell'apprendimento:

• Decomporre il formato floating-point per calcolare valori numerici da schemi binari usando la rappresentazione normalizzata e la codifica in eccesso.
• Visualizzare la distribuzione dei numeri rappresentabili sulla linea numerica e spiegare l'impatto della suddivisione dei bit tra esponente e mantissa.
• Quantificare l'inaccuratezza numerica usando ULP e identificare come diversi modi di arrotondamento contribuiscono all'errore.

🔹 Lezione 8: Studio di caso: Parallelizzazione della ricostruzione MRI

Panoramica: Questa lezione esplora la parallelizzazione della ricostruzione avanzata di Immagini a Risonanza Magnetica (MRI) su GPU. Si concentra sul processo iterativo di ricostruzione per traiettorie non cartesiane, ottimizzando in particolare il kernel computazionalmente intensivo F^H d tramite trasformazioni del ciclo, gestione della memoria costante, riorganizzazione della disposizione dei dati e utilizzo di funzioni trigonometriche accelerate a livello hardware.

Risultati dell'apprendimento:

• Comprendere il passaggio dalla ricostruzione basata su FFT cartesiana a algoritmi iterativi basati su risolutore lineare per dati k-space non cartesiani.
• Applicare la divisione del ciclo (loop fission) e lo scambio di ciclo (loop interchange) per trasformare codice C sequenziale in una struttura adatta a una mappatura massiva di thread CUDA.
• Ottimizzare il throughput della memoria usando chunking della memoria costante e layout dei dati Array-of-Structs (AoS).

🔹 Lezione 9: Studio di caso: Visualizzazione molecolare ed esecuzione multi-GPU

Panoramica: Questa lezione esplora l'applicazione pratica del calcolo GPU alla visualizzazione molecolare, specificamente utilizzando il metodo della Somma Coulomb Diretta (DCS) per calcolare mappe del potenziale elettrostatico. Gli studenti passeranno da un'implementazione di kernel di base a versioni altamente ottimizzate che sfruttano l'unrolling delle istruzioni, il coalescing della memoria e il padding.

Risultati dell'apprendimento:

• Implementare un kernel Somma Coulomb Diretta (DCS) utilizzando la memoria costante CUDA e tecniche di latenza nascosta della memoria globale.
• Ottimizzare le prestazioni del kernel tramite unrolling delle istruzioni e il riutilizzo di calcoli comuni delle coordinate.
• Applicare strategie di coalescing e padding per allineare gli accessi alla memoria globale GPU al fine di massimizzare la larghezza di banda.

🔹 Lezione 10: Pensiero computazionale e selezione di algoritmi paralleli

Panoramica: Questa lezione esplora il passaggio dal pensiero sequenziale a quello parallelo focalizzandosi sugli obiettivi della programmazione parallela e sulla scelta strategica di algoritmi. Gli studenti impareranno a decomporre problemi in unità parallelizzabili, applicare il pensiero computazionale per colmare il divario tra scienza del dominio e architettura hardware, e valutare le prestazioni degli algoritmi.

Risultati dell'apprendimento:

• Identificare gli obiettivi principali della programmazione parallela e calcolare il guadagno teorico in velocità usando la legge di Amdahl.
• Differenziare tra decomposizione a livello di task e a livello di dati e applicare strategie atom-centriche (scatter) rispetto a grid-centriche (gather).
• Valutare e selezionare algoritmi paralleli in base a criteri come larghezza di banda della memoria, complessità computazionale e vincoli architetturali.

🔹 Lezione 11: Introduzione al modello di programmazione OpenCL

Panoramica: Questa lezione introduce OpenCL come framework per il calcolo parallelo eterogeneo, con particolare enfasi sul suo modello di parallelismo dati e sull'astrazione gerarchica dell'hardware. Gli studenti impareranno a mappare le strutture NDRange e di memoria di OpenCL su equivalenti CUDA e padroneggiare la gestione lato host dei dispositivi tramite un modello di compilazione dinamica.

Risultati dell'apprendimento:

• Mappare il parallelismo e le gerarchie di memoria di OpenCL su architetture specifiche CUDA (es. mappare Work-groups a Blocks e Local Memory a Shared Memory).
• Implementare funzioni kernel OpenCL e gestire l'ambiente di esecuzione lato host tramite Context e Command Queues.
• Eseguire il flusso di lavoro di compilazione dinamica per creare kernel da codice sorgente in tempo reale.

🔹 Lezione 12: Funzionalità moderne delle GPU e prospettive future

Panoramica: Questa lezione esplora l'evoluzione architetturale e funzionale delle GPU, concentrando l'attenzione sul passaggio verso una gestione della memoria più sofisticata, capacità migliorata di esecuzione dei kernel e aumento della performance dei core. Gli studenti esamineranno come funzionalità come lo Spazio di Memoria Unificato del Dispositivo e le chiamate di funzione a livello di kernel trasformino la GPU in un processore general-purpose.

Risultati dell'apprendimento:

• Spiegare l'importanza dell'Evoluzione dell'Architettura della Memoria e il passaggio verso uno Spazio di Memoria Unificato a 64 bit.
• Analizzare come Operazioni Atomiche Potenziate e Chiamate di Funzione a Livello di Kernel consentono l'implementazione di strutture e algoritmi complessi.
• Valutare gli impatti sulle prestazioni dell'Esecuzione Simultanea di Kernel, degli Miglioramenti nella Velocità a Precisione Doppia e dell'Efficienza del Flusso di Controllo in ambienti GPU moderni.