Corso di formazione Programmazione C++ accelerata con CUDA

Introduzione al calcolo accelerato con le GPU

• Il concetto di elaborazione eterogenea e l’architettura CPU-GPU
• Esecuzione del codice CUDA e gestione degli spazi di memoria
• Compilazione del codice C++ per GPU tramite nvcc e CMake
• Verifica dell’ambiente di sviluppo relativo alla GPU

Algoritmi paralleli con Thrust e CUB

• Esecuzione su GPU di operazioni come ordinamento, riduzione e trasformazione dei dati
• Ristrutturazione di algoritmi della STL affinché possano essere eseguiti sulla GPU
• Utilizzo dei vettori dispositivi di Thrust e delle relative strategie di esecuzione
• Utilizzo dei primitivi forniti da CUB per costruire pipeline parallele personalizzate

Architettura e gestione della memoria GPU

• I tipi di memoria: globale, costante e da texture
• Assegnazione e trasferimento esplicito della memoria tra CPU e GPU
• Memoria Unificata per una gestione più semplice dell’accesso ai dati
• Ottimizzazione delle modalità di accesso alla memoria tramite allineamento e raggruppamenti logici

Esecuzione asincrona mediante gli stream CUDA

• Creazione e gestione degli stream di esecuzione CUDA
• Sovrapposizione tra l’esecuzione dei kernel e i trasferimenti dati in memoria
• Utilizzo degli eventi CUDA per gestire le dipendenze temporali
• Ottimizzazione della concorrenza tramite impostazioni di priorità degli stream

Sviluppo dei propri kernel CUDA

• Comprensione del modello di programmazione SIMT e dell’esecuzione per warp
• Configurazione del lancio dei kernel, inclusi i cicli a dimensione variabile
• Gestione degli indici thread e delle griglie multidimensionali
• Rilevamento ed elaborazione degli errori tramite l’API runtime di CUDA

Gerarchia dei thread e modello d’esecuzione

• Relazioni tra griglie, blocchi e singoli thread nel codice dispositivo
• Operazioni atomiche a livello di warp e meccanismi di voto interno ai gruppi
• Sincronizzazione dei thread all’interno degli stessi blocchi
• Analisi dell’occupancy e delle risorse hardware impiegate durante l’esecuzione

Gruppi cooperativi per un parallelismo flessibile

• L’API cooperative_groups e i diversi tipi di gruppi che si possono definire
• Implementazione basata su tasselli (tiles) con il metodo tile_partition
• Lancio di operazioni parallele a livello di intera griglia
• Schema sincronizzazione multi-griglia per applicazioni complesse

Tecniche d’ottimizzazione con la memoria condivisa

• Evitare i conflitti di indirizzamento nella memoria condivisa
• Sfruttare strategie di suddivisione dei dati per operazioni matriciali
• Utilizzo della memoria condivisa come cache gestita direttamente dall’utente
• Uso della classe cuda::shared_memory_mdspan per rappresentare array multidimensionali su memoria condivisa

Fusione dei kernel e pattern paralleli avanzati

• Unione di più kernel in uno solo al fine di ridurre il overhead d’esecuzione
• Implementazione di algoritmi come la scansione, l’aggregazione per chiavi e operazioni segmentate
• Uso delle operazioni atomiche senza blocchi e strutture dati thread-safe
• Utilizzo dei meccanismi atomici aggregati a livello di warp per aumentare la velocità d’elaborazione

Analisi prestazionale con Nsight Systems

• Visualizzazione temporale delle attività CPU e GPU
• Rilevamento dei colli di bottiglia nei trasferimenti di dati tra memoria centrale e memoria grafica
• Analisi del rendimento dei kernel in termini di utilizzo delle risorse hardware
• Ottimizzazione iterativa con l’utilizzo di Nsight Compute per ulteriori approfondimenti

Utilizzo avanzato delle feature moderne del linguaggio C++ nei kernel CUDA

• Impiego di lambda expressions, costanti compile-time (constexpr) e tipi dinamici (auto) all’interno dei kernel
• Riconoscimento e utilizzo degli algoritmi paralleli introdotti in C++17
• Sfruttamento delle espressioni logiche definite nei concetti di C++20 e la gestione degli intervalli di dati (Ranges) sul dispositivo
• Supporto alle funzionalità introdotte in C++23 grazie a nvcc e alla libreria CCCL 3.x

Grafici CUDA ed esecuzione asincrona evoluta

• Definizione e lancio di grafici d’esecuzione basati su CUDA
• Cattura automatica dei comandi eseguiti tramite gli stream per generare tali grafici
• Modifica dinamica dei nodi contenuti nei grafici e utilizzo di condizioni operative
• Riduzione dell’overhead di lancio nei processi iterativi e complessi

Integrazione delle librerie CUDA nelle applicazioni preesistenti

• Sviluppo di interfacce C++ utilizzabili per incapsulare il codice GPU
• Gestione ottimale dei sistemi dotati di più GPU o configurazioni NUMA
• Integrazione nel flusso di compilazione mediante strumenti come CMake e CUDA
• Risoluzione dei problemi logici nel codice dispositivo con l’utilizzo di cuda-gdb

Considerazioni finali e best practice

• Scegliere quando adottare Thrust, CUB o kernel personalizzati per specifiche necessità applicative
• Assicurare la portabilità delle soluzioni sviluppate su diverse tipologie di GPU
• Gestione strutturata del codice e uso dei principi RAII per la gestione delle risorse GPU
• Suggerimenti sui prossimi passi da compiere per approfondire le conoscenze su CUDA