Aller au contenu

PGPU35 · Programmation GPU

L'UE dont dépendent vos chiffres de performance en compétition. La majorité des machines du Top500 et du Green500 sont hétérogènes : sans GPU, on ne joue plus.

Fiche signalétique
Code PGPU35
Crédits 5 ECTS
Semestre 5, spécialisation — mardi après-midi
Responsable ROUSSEL Adrien
Concurrentes sur le créneau PRRD35, INCA35, MALE35
Choix Une des trois UE de spécialisation, en accord avec l'entreprise
Prérequis Programmation parallèle ; C/C++
Effectif max 30

Ce que dit la brochure

Objectifs cités :

« La programmation sur les cartes graphiques (GPGPU : General Purpose programming on GPU) a commencé en 2008 avec l'avènement de l'architecture CUDA (Compute Unified Device Architecture) par Nvidia. Aujourd'hui la majorité des supercalculateurs du Top500 et du Green500 sont basées sur des architectures hétérogènes (CPU+accélérateurs GPU). Les cartes graphiques se retrouvent également dans de nombreuses machines, de l'embarqué, en passant par les smartphones, les machines de bureau, stations de travail, etc. L'objectif de cette UE est d'exploiter efficacement ces plateformes hétérogènes exposants des processeurs classiques CPUS et des accélérateurs GPUs. »

Contenu cité : « L'architecture CUDA de Nvidia sera utilisée, mais les concepts sont applicables aux autres architectures HIP d'AMD ou SYCL. »

  • écosystème CUDA : domaines concernés, comparaison avec le CPU, matériels et logiciels (compilateur, profilage, bibliothèques) ;
  • architecture des GPU : organisation des unités de calcul, hiérarchie mémoire, interaction avec la machine hôte ;
  • modèle de programmation : modèle SPMD, notion de kernel, threads, groupes de threads, synchronisation et communication entre threads ;
  • modèle d'exécution : placement de blocs de threads sur des blocs de cœurs, contraintes matérielles ;
  • programmation : API, fonctions de base, mesure du temps des opérations GPU, gestion des erreurs, erreurs courantes ;
  • portage de codes C/C++ vers GPU ;
  • optimisation : dimensionnement des grilles et blocs de threads, allocation mémoire, fonctionnement asynchrone, recouvrement des calculs et des communications, accès à la mémoire globale, mémoire shared ;
  • autres points : multi-GPU, MPI+GPU, utilisation de Tensor Cores.

Ce que ça vaut pour le HPC

Maximal. C'est, avec PRPA23 et ICPA24, la troisième jambe du parallélisme, et celle qui pèse le plus lourd dans les chiffres bruts.

Quelques ordres de grandeur pour situer l'enjeu. Un processeur serveur récent délivre typiquement quelques téraflops en double précision ; un accélérateur de calcul récent en délivre plusieurs dizaines, avec une bande passante mémoire cinq à dix fois supérieure. Sur un benchmark comme HPL, plus de 90 % de la performance d'un nœud hétérogène vient des accélérateurs.

Trois points du programme méritent une attention particulière.

1. Le recouvrement des calculs et des communications. C'est la compétence qui distingue un code GPU médiocre d'un code GPU efficace. Le transfert de données entre l'hôte et l'accélérateur passe par un bus dont la bande passante est très inférieure à celle de la mémoire du GPU. Un code naïf passe la majorité de son temps à transférer. La solution — flux CUDA (streams), transferts asynchrones, découpage du travail en tranches recouvertes — est au programme.

2. Multi-GPU et MPI+GPU. Un nœud de calcul moderne porte quatre à huit accélérateurs, reliés entre eux par un réseau interne rapide. Exploiter plusieurs GPU sur un nœud, puis plusieurs nœuds, suppose de combiner CUDA et MPI. Le GPU-aware MPI permet de passer directement un pointeur de mémoire d'accélérateur à MPI_Send, la bibliothèque prenant en charge le chemin le plus court — potentiellement sans passer par la mémoire de l'hôte. C'est un savoir opérationnel rare.

3. Tensor Cores. Les unités matricielles dédiées, initialement pour l'IA, sont aussi exploitables en calcul scientifique via la précision mixte. C'est ce que mesure le benchmark HPL-MxP, au programme de la compétition SC26.

Ressource interne : un cours complet déjà disponible

Cette bibliothèque contient un cours progressif et complet sur la programmation GPU : Développement GPU : de zéro à expert, qui va de « qu'est-ce qu'un warp » jusqu'aux techniques récentes.

Lisez-le pendant l'été précédant le semestre 5. Vous arriverez au premier cours avec l'architecture et le modèle de programmation déjà acquis, et vous pourrez consacrer l'UE entière à l'optimisation — qui est la partie difficile.

Arriver prêt

⏱ 25 h, et c'est l'investissement le plus rentable du semestre.

  1. Lire le cours GPU de cette bibliothèque (lien ci-dessus). ⏱ 15 h.
  2. Installer CUDA et faire tourner un premier kernel. Si vous n'avez pas de carte NVIDIA, les options sont : Google Colab (gratuit, accès GPU limité), une instance cloud à la demande, ou l'accès à une machine de l'école ou d'un centre. ⏱ 3 h.
  3. Écrire et mesurer trois kernels avant le cours : une addition de vecteurs (SAXPY), une réduction, et une transposition de matrice. Mesurer la bande passante atteinte et la comparer à la bande passante théorique du GPU. La transposition naïve atteindra environ 10 % du maximum : comprendre pourquoi (accès non coalescés) est le point d'entrée de toute l'optimisation GPU. ⏱ 5 h.
  4. Prendre en main Nsight Systems et Nsight Compute. ⏱ 2 h. Le profilage GPU est plus déterminant que sur CPU, parce que l'intuition y est encore moins fiable.

Ressources

Priorité 1

  • Wen-mei Hwu, David Kirk, Izzat El Hajj, Programming Massively Parallel Processors: A Hands-on Approach, 4ᵉ éd., Morgan Kaufmann, 2022. Le livre de référence, écrit par les auteurs du premier cours universitaire de CUDA. La 4ᵉ édition couvre les architectures récentes et les Tensor Cores.
  • Le CUDA C++ Programming Guide de NVIDIA. La documentation officielle, et elle est bonne. À lire au moins une fois en entier : les chapitres sur le modèle mémoire et sur les streams sont incontournables.
  • Le CUDA C++ Best Practices Guide, du même éditeur. Plus court, organisé par priorité d'optimisation. C'est la liste de contrôle à suivre.
  • Le blog technique de NVIDIA (developer.nvidia.com/blog), séries « CUDA Pro Tip » et « How to Optimize Data Transfers », « How to Overlap Data Transfers ». Courts, précis, avec du code.

Priorité 2 — l'optimisation avancée

  • Les exposés de la GTC (GPU Technology Conference), archivés et gratuits. Chercher notamment les sessions sur la microarchitecture, sur le profilage avec Nsight, et les exposés historiques de Mark Harris sur les réductions (Optimizing Parallel Reduction in CUDA), qui restent le meilleur tutoriel d'optimisation GPU jamais écrit — sept versions successives d'une réduction, chacune mesurée.
  • La documentation de Nsight Compute : les métriques, les sections d'analyse, et le modèle d'occupation. Comprendre la notion d'occupancy et savoir qu'une occupancy élevée n'est pas toujours souhaitable.
  • Les bibliothèques : cuBLAS, cuSPARSE, cuFFT, cuSOLVER, cuDNN, Thrust, CUB, NCCL. Règle d'or : ne réécrivez jamais ce que cuBLAS fait déjà, sauf comme exercice. Savoir appeler la bibliothèque correctement est la compétence utile.
  • Sur le multi-GPU et MPI : la documentation de NCCL pour les collectives entre accélérateurs, et les pages d'OpenMPI et de MPICH sur le support CUDA (CUDA-aware MPI).

Priorité 3 — la portabilité

La brochure le dit lui-même : les concepts sont transposables. Et c'est stratégique, parce que les machines récentes ne sont pas toutes NVIDIA.

  • HIP (AMD) : presque une traduction mot à mot de CUDA. L'outil hipify convertit automatiquement la majeure partie d'un code. Important : plusieurs grandes machines européennes et américaines sont à accélérateurs AMD.
  • SYCL (standard Khronos) et son implémentation oneAPI DPC++ d'Intel, ou AdaptiveCpp. Du C++ standard, portable sur CPU, GPU NVIDIA, GPU AMD et GPU Intel.
  • Kokkos et RAJA : les couches d'abstraction utilisées par les codes exascale américains. Lien direct avec LAOA23.
  • OpenMP target : le chemin de portage par directives, vu en ICPA24, et OpenACC, son concurrent historique.
  • nvc++ -stdpar : exécute les algorithmes parallèles du C++17 sur GPU sans écrire une ligne de CUDA. Spectaculaire et peu connu.

Outils à maîtriser

Outil Usage
nvcc, nvc++ Compilation ; comprendre -arch=sm_XX et la compilation JIT
nvidia-smi État, mémoire, utilisation, température, puissance électrique
nvidia-smi dmon, nvidia-smi -q -d POWER Suivi de la consommation — utile en compétition
Nsight Systems (nsys) Chronologie : voir les trous où le GPU attend
Nsight Compute (ncu) Analyse d'un kernel : occupancy, débits, roofline
compute-sanitizer L'équivalent des sanitizers pour CUDA : accès mémoire, courses
cuda-gdb Débogage
CUDA events Chronométrage correct côté GPU
rocm-smi, rocprof Les équivalents AMD

L'erreur de mesure numéro un sur GPU

Les appels CUDA sont asynchrones. Ce code mesure le temps de soumission, pas d'exécution, et donnera un résultat absurdement bon :

t0 = now();
mon_kernel<<<grille, bloc>>>(...);
t1 = now();            // FAUX : le kernel n'est probablement pas terminé

La version correcte :

cudaEvent_t debut, fin;
cudaEventCreate(&debut); cudaEventCreate(&fin);
cudaEventRecord(debut);
mon_kernel<<<grille, bloc>>>(...);
cudaEventRecord(fin);
cudaEventSynchronize(fin);
float ms; cudaEventElapsedTime(&ms, debut, fin);

Et toujours vérifier les codes de retour. Un kernel qui échoue ne signale rien : il ne s'exécute simplement pas, et votre mesure est celle d'une opération vide. Enveloppez tous les appels dans une macro de vérification, et appelez cudaGetLastError() après chaque lancement.

Exercices

E1 · ★ ⏱ 2 h — SAXPY et la bande passante. Écrire \(y \leftarrow \alpha x + y\) en CUDA pour un vecteur de 100 millions d'éléments. Calculer la bande passante effective : trois accès mémoire (deux lectures, une écriture) de 4 octets par élément, divisés par le temps. Comparer à la bande passante théorique de la carte (sur la fiche technique) et à celle mesurée par un banc d'essai comme BabelStream. Vous devriez atteindre 70 à 90 %. Si vous êtes à 20 %, quelque chose ne va pas.

E2 · ★★ ⏱ 3 h — La coalescence. Une transposition de matrice \(4096^2\), en trois versions : naïve (lecture coalescée, écriture dispersée), avec mémoire shared comme tampon de bloc, et avec mémoire shared plus remplissage (padding) pour éviter les conflits de banc. Mesurer les trois. L'écart entre la première et la dernière est typiquement d'un facteur cinq à dix. C'est l'exercice le plus instructif de tout CUDA : il enseigne la coalescence, la mémoire partagée et les conflits de bancs en une seule fois.

E3 · ★★ ⏱ 3 h — La réduction, sept fois. Reproduire les sept versions successives de la réduction du tutoriel de Mark Harris : entrelacée avec divergence, entrelacée sans divergence, séquentielle, premier niveau déroulé, dernier warp déroulé, complètement déroulée par template, et plusieurs éléments par thread. Mesurer chacune. Puis comparer à thrust::reduce et à cub::Reduce. C'est l'exercice qui apprend à penser en warps.

E4 · ★★★ ⏱ 4 h — Recouvrir les transferts. Un calcul sur un tableau qui ne tient pas dans la mémoire du GPU : le traiter par tranches. Version 1 : transfert, calcul, transfert retour, en série pour chaque tranche. Version 2 : plusieurs streams CUDA, avec transferts asynchrones et mémoire hôte épinglée (pinned), de sorte que le transfert de la tranche \(i+1\) recouvre le calcul de la tranche \(i\). Mesurer, et visualiser la chronologie dans Nsight Systems : vous verrez littéralement les barres se superposer. Le gain théorique maximal est un facteur deux à trois.

E5 · ★★★ ⏱ 5 h — GEMM à la main. Écrire un produit matrice-matrice en CUDA, par étapes : version naïve un thread par élément, version avec tiling en mémoire partagée, version avec plusieurs éléments par thread et déroulement. Mesurer en GFLOPS et comparer à cuBLAS. Objectif réaliste : atteindre 40 à 60 % de cuBLAS. Atteindre 90 % est un travail de plusieurs semaines — et c'est ce que fait cuBLAS. La leçon : on n'écrit pas son GEMM, on comprend pourquoi.

E6 · ★★★ ⏱ 5 h — Multi-GPU. Distribuer un calcul sur tous les GPU d'un nœud, de deux manières : avec l'API CUDA multi-appareil (cudaSetDevice) et avec un rang MPI par GPU. Mesurer l'accélération et le coût de la synchronisation. Puis tester le GPU-aware MPI : passer directement des pointeurs d'appareil à MPI_Allreduce et comparer à la version qui recopie par l'hôte.

E7 · ★★★★ ⏱ 6 h — Précision mixte et Tensor Cores. Implémenter un produit matriciel utilisant les Tensor Cores, via cublasGemmEx en précision mixte (entrées FP16 ou BF16, accumulation FP32), et comparer aux GFLOPS de la version FP32 et FP64. Puis étudier l'effet sur la précision du résultat : mesurer l'erreur relative par rapport à une référence FP64. Enfin, implémenter un raffinement itératif — résoudre en précision réduite puis corriger en précision supérieure — et constater qu'on peut retrouver la précision FP64 à une vitesse proche du FP16. C'est exactement le principe du benchmark HPL-MxP.

Projet

Projet PGPU35 · Le portage complet d'un code sur GPU

⏱ 40 h · ★★★★

Sujet. Prendre un code CPU existant, non trivial, et le porter sur GPU en documentant chaque décision et chaque mesure.

Choix du code. Idéalement, le vôtre : le solveur de Poisson de PRPA23, ou le noyau optimisé de ICPA24. Sinon, un mini-application reconnue : les proxy apps du projet Exascale américain — LULESH (hydrodynamique), miniFE (éléments finis), XSBench (transport de neutrons), BabelStream (bande passante) — sont conçues pour cet usage, petites et représentatives.

Les huit étapes :

  1. Mesure de référence CPU, optimisée et parallélisée. Il ne s'agit pas de comparer un GPU réglé à un CPU volontairement mauvais : c'est la faute de méthode la plus fréquente dans la littérature GPU.
  2. Modèle de performance a priori : intensité arithmétique du noyau, roofline du GPU, prédiction du régime et de la performance atteignable. Écrire la prédiction avant d'implémenter.
  3. Portage naïf, correct, validé numériquement contre la version CPU avec une tolérance justifiée.
  4. Profilage avec Nsight Compute : occupancy, débit mémoire, efficacité des accès, divergence de warps. Confronter au modèle.
  5. Optimisations successives, mesurées une par une : coalescence, mémoire partagée, dimensionnement des blocs, déroulement, réduction des transferts.
  6. Recouvrement calcul-transfert avec des streams.
  7. Passage à plusieurs GPU, avec courbe d'accélération.
  8. Comparaison avec une voie alternative : refaire le portage avec OpenMP target, ou avec Kokkos, ou avec SYCL, et comparer performance obtenue et effort de développement. C'est la question que se posent réellement les équipes qui portent de gros codes, et personne ne la documente honnêtement.

Livrables : le code, un tableau des optimisations avec gains mesurés, un graphique roofline avec la trajectoire de vos versions, la chronologie Nsight avant et après recouvrement, et un rapport de dix pages.

Pourquoi ce projet. Parce que le portage de codes sur accélérateur est, en 2026, l'activité la plus demandée du HPC. Tous les grands codes de simulation y passent, et il manque des gens capables de le faire proprement — c'est-à-dire avec un modèle, des mesures et une validation numérique.

Erreurs fréquentes

Neuf pièges du GPU

  1. Mesurer sans synchroniser. Voir l'encadré ci-dessus. L'erreur la plus fréquente et la plus grossière.
  2. Ne pas vérifier les erreurs CUDA. Un kernel qui ne s'exécute pas ne dit rien. Toujours une macro de vérification.
  3. Oublier le coût du transfert. Un kernel dix fois plus rapide que le CPU, précédé et suivi de transferts qui coûtent vingt fois le calcul, est une régression. Mesurez toujours le temps de bout en bout, transferts inclus.
  4. Accès non coalescés. Si les threads consécutifs d'un warp n'accèdent pas à des adresses consécutives, la bande passante s'effondre. C'est la première chose à vérifier.
  5. Divergence de warps. Un if dont le résultat diffère entre threads d'un même warp force l'exécution des deux branches en série. Dans une boucle chaude, c'est coûteux.
  6. Croire qu'il faut maximiser l'occupancy. Une occupancy de 100 % n'est pas un objectif : un kernel qui utilise plus de registres par thread a une occupancy plus faible mais peut être plus rapide. La métrique est le débit, pas l'occupancy.
  7. Comparer à un CPU mal optimisé. « 100× plus rapide que le CPU » signifie presque toujours « comparé à une version séquentielle non vectorisée compilée en -O0 ». Les comparaisons honnêtes donnent des facteurs de 3 à 10 par rapport à un CPU bien optimisé, ce qui est déjà considérable.
  8. Réécrire ce que cuBLAS fait. Faites-le une fois comme exercice, jamais en production.
  9. Ignorer la mémoire épinglée. Les transferts depuis de la mémoire hôte ordinaire sont deux fois plus lents que depuis de la mémoire épinglée, et ils ne peuvent pas être asynchrones. cudaHostAlloc ou cudaMallocHost : une ligne, un facteur deux.

Comment l'UE s'articule avec le reste

UE ou chapitre Lien
LAOA23 (S3) CUDA est du C++ ; Thrust est la STL sur GPU
PRPA23 (S3) MPI+GPU, et le GPU-aware MPI
ICPA24 (S4) L'offload OpenMP est la voie alternative
MALE24 (S4) L'apprentissage tourne sur GPU ; Tensor Cores
COAV35 (S5) La compilation pour accélérateur
GIIG35 (S5) Le Green500 est dominé par les architectures hétérogènes
Cours GPU de cette bibliothèque Le cours complet, à lire avant
GPU et hétérogène Le complément : NCCL, entraînement distribué, HIP, SYCL

À retenir

PGPU35 en trois phrases

C'est l'UE qui décide de vos chiffres bruts : la majorité du Top500 et du Green500 est hétérogène, et sur un nœud à accélérateurs, plus de 90 % de la performance vient du GPU. Les trois compétences à maîtriser absolument sont la coalescence des accès mémoire, le recouvrement calcul-transfert par streams, et la mesure correcte avec des CUDA events. Lisez le cours GPU de cette bibliothèque pendant l'été : vous consacrerez alors l'UE entière à l'optimisation, qui est la partie difficile.

Fiche suivante : PDSP35 · Performance des systèmes parallèles.