3 · Papiers de recherche¶
Classés par thème, avec pour chacun ce qu'il apporte et quand le lire.
3.1 Les fondations¶
Modèle de performance¶
Williams, Waterman, Patterson — Roofline: An Insightful Visual Performance Model for Multicore Architectures, CACM 52(4), 2009. DOI 10.1145/1498765.1498785
Le papier qui a donné son nom au modèle. Court, lisible, et son idée centrale n'a pas bougé en dix-sept ans. À lire avant la partie 3.
Yang et al. — Hierarchical Roofline Performance Analysis for Deep Learning Applications. arXiv:2009.05257
L'extension à plusieurs niveaux de mémoire, telle qu'implémentée dans Nsight Compute.
Programmation persistante¶
Gupta, Stuart, Owens — A Study of Persistent Threads Style GPU Programming for GPGPU Workloads, InPar 2012. eScholarship · présentation GTC 2012
Le papier fondateur des megakernels, quatorze ans avant qu'ils portent ce nom. Il définit formellement le style persistent threads, identifie ses quatre cas d'usage, et conclut honnêtement qu'il « peut atteindre jusqu'à un ordre de grandeur d'accélération mais peut aussi entraîner une perte de performance dans de nombreux cas ».
À lire avant la partie 8.
Sorensen et al. — Cooperative Kernels: GPU Multitasking for Blocking Algorithms. arXiv:1707.01989
Sur les garanties de progression et le multitâche des noyaux coopératifs.
Occupancy¶
Volkov — Better Performance at Lower Occupancy, GTC 2010. PDF
Une présentation, pas un papier, mais elle a plus influencé la pratique que la plupart des papiers. Elle démontre que le parallélisme d'instructions substitue au parallélisme de threads, et qu'une occupancy faible peut être optimale.
Nicolas et al. — Optimization Techniques for GPU Programming, ACM Computing Surveys 55(11), 2023. ACM DL
Un état de l'art synthétique des techniques d'optimisation. Bon point d'entrée bibliographique.
3.2 La microarchitecture¶
Ces papiers démontent le matériel instruction par instruction. Ils sont la meilleure source de chiffres réels, souvent absents des documentations.
| Papier | Cible | Ce qu'on y trouve |
|---|---|---|
| Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking, arXiv:1804.06826 | Volta | tensor cores 1re gén., latences |
| Dissecting the NVidia Turing T4 GPU via Microbenchmarking, arXiv:1903.07486 | Turing | hiérarchie mémoire |
| Dissecting the NVIDIA Hopper Architecture, arXiv:2501.12084 | Hopper | wgmma à 95 % du pic contre 62,9 % pour mma, TMA, clusters |
| Microbenchmarking NVIDIA's Blackwell Architecture, arXiv:2512.02189 | Blackwell | latences tcgen05 (~11 cycles), TMEM, moteur de décompression |
Pourquoi ces papiers comptent
NVIDIA ne publie pas les latences d'instructions, les tailles de files, ni le comportement exact des caches. Ces papiers les mesurent.
Le papier Blackwell donne par exemple les débits par précision (FP16 :
1 929 TFLOPS, FP8 : 3 851, FP4 : 7 702) et les latences de tcgen05.mma
(11,0 à 11,4 cycles selon la tuile), chiffres qu'on ne trouve nulle part
ailleurs.
3.3 Les noyaux d'IA¶
Attention¶
Dao, Fu, Ermon, Rudra, Ré — FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness, NeurIPS 2022. arXiv:2205.14135
Le papier le plus influent de la programmation GPU pour l'IA. Pavage, softmax en ligne, recalcul. À lire intégralement.
Dao — FlashAttention-2. arXiv:2307.08691
Réduction des opérations non-matmul, meilleure répartition du travail.
Shah et al. — FlashAttention-3. arXiv:2407.08608
Exploitation de Hopper : TMA, wgmma, chevauchement softmax/GEMM, FP8.
— FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling, mars 2026. arXiv:2603.05451
Le co-design pour Blackwell : exponentielles émulées sur les unités FMA, remise à l'échelle conditionnelle (~10× moins), écrit entièrement en CuTe DSL. 1 605 TFLOPS sur B200, 1,3× cuDNN, 2,7× Triton.
Voir aussi la rétro-ingénierie par Modal.
Quantification¶
| Papier | Apport |
|---|---|
| GPTQ, arXiv:2210.17323 | quantification post-entraînement avec compensation hessienne |
| AWQ, arXiv:2306.00978 | protection des canaux importants, identifiés par les activations |
| SmoothQuant, arXiv:2211.10438 | déplacement de la difficulté des activations vers les poids |
| OCP Microscaling Formats (MX) | la spécification des formats à échelle par blocs |
Service et inférence¶
Kwon et al. — Efficient Memory Management for Large Language Model Serving with PagedAttention (vLLM), SOSP 2023. arXiv:2309.06180
Le papier de vLLM. La gestion du cache KV par pages, inspirée de la mémoire virtuelle.
Leviathan, Kalman, Matias — Fast Inference from Transformers via Speculative Decoding, ICML 2023. arXiv:2211.17192
Cai et al. — Medusa. arXiv:2401.10774
— Memory-Bound but Not Bandwidth-Limited: The Physical AI Inference Gap in Batch-1 LLM Decode, mai 2026. arXiv:2605.30571
Une analyse du régime de décodage à lot 1 qui complète directement la partie 7.
Noyaux et compilateurs¶
Tillet, Kung, Cox — Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations, MAPL 2019. ACM DL
Le papier fondateur de Triton.
— Liger Kernel: Efficient Triton Kernels for LLM Training. arXiv:2410.10989
— The Anatomy of a Triton Attention Kernel. arXiv:2511.11581
— Optimal Software Pipelining and Warp Specialization for Tensor Core GPUs. arXiv:2512.18134
3.4 Le corpus des megakernels¶
C'est le cœur de la partie 8. Dans l'ordre chronologique.
2012 — la préhistoire¶
Gupta, Stuart, Owens, Persistent Threads (voir §3.1).
Mai 2025 — le premier megakernel de modèle complet¶
Hazy Research — Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B. Blog · code · discussion HN
Un billet de blog, pas un papier — et le document fondateur du domaine. Interpréteur sur GPU, allocateur de pages, compteurs. Llama-1B en moins d'une milliseconde sur H100, 78 % de la bande passante.
Septembre 2025 — la variante débit¶
Hazy Research — We Bought the Whole GPU, So We're Damn Well Going to Use the Whole GPU. Blog
Llama-70B en tensor-parallèle sur 8 H100. Recouvrement à trois niveaux, transposition distribuée, +22 % sur SGLang.
2025 — l'approche AMD¶
Kog — Building a single-kernel, latency-optimized LLM inference engine on AMD MI300X GPUs. Blog
Sentinelles NaN (sync 7,6 → 0,9 µs), duplication par XCD, streaming continu des poids, parallélisme de tenseurs différé.
Décembre 2025 — le compilateur¶
Cheng et al. — Mirage Persistent Kernel: A Compiler and Runtime for Mega-Kernelizing Tensor Programs. arXiv:2512.22219 · code · page projet
Le papier de référence du domaine. ttGraph, événements, workers et schedulers, lancement hybride JIT/AOT. 1,0-1,7× contre SGLang/vLLM.
Avril 2026 — le dynamisme¶
Jin et al. — Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel. arXiv:2604.13327
Formes et dépendances dynamiques. Le verrou qui décidera de l'avenir de la technique.
Avril 2026 — les chiplets¶
Chowdhary et al. (AMD) — Fleet: Hierarchical Task-based Abstraction for Megakernels on Multi-Die GPUs. arXiv:2604.15379
Chiplet-tasks. Sur MI350/Qwen3-8B : 1,3-1,5× vs vLLM, L2 de 12 % à 54 %, −37 % de trafic HBM.
Mai 2026 — la production¶
Dong et al. (Baidu) — Ada-MK: Adaptive MegaKernel Optimization via Automated DAG-based Search for LLM Inference. arXiv:2605.11581
Recherche DAG hors ligne, MLIR, découpage en K. Déployé en production dans le système publicitaire de Baidu. +23,6 % vs TensorRT-LLM, +50,2 % vs vLLM.
Juin 2026 — la synthèse automatique¶
Jaber & Jaber — AutoMegaKernel: A Statically-Checked Agent Harness for Self-Retargeting Megakernel Synthesis. arXiv:2606.09682
Validateur d'IR figé, agent proposant des ordonnancements, rejet avant exécution. 7 160 ordonnancements adverses validés, zéro faux positif, parité jeton à jeton, 1,25-1,72×.
Travaux connexes¶
| Papier | Sujet |
|---|---|
| SonicMoE, arXiv:2512.14080 | MoE, optimisations E/S et tuiles |
| Deep Kernel Fusion for Transformers, arXiv:2602.11808 | fusion profonde |
| Tile-Level Activation Overlap, arXiv:2607.02521 | recouvrement au niveau des tuiles |
| TokenWeave, arXiv:2505.11329 | recouvrement calcul/communication |
3.5 Communication et multi-GPU¶
Shoeybi et al. — Megatron-LM: Training Multi-Billion Parameter Language Models Using Model Parallelism. arXiv:1909.08053
Le parallélisme de tenseurs colonnes/lignes.
— Demystifying NVSHMEM: A System-Level Analysis on Symmetric Memory and Device-Initiated Operations in GPU Communication, juin 2026. arXiv:2606.05951
L'analyse de référence sur NVSHMEM, la mémoire symétrique et les opérations initiées par le périphérique. Couvre aussi DeepEP en détail.
— GPU-Initiated Networking for NCCL. arXiv:2511.15076
3.6 Génération automatique de noyaux¶
Un domaine très actif en 2026.
| Papier | Apport |
|---|---|
| KernelBench-X, arXiv:2605.04956 | le benchmark de référence, étendu |
| FlashInfer-Bench, arXiv:2601.00227 | évaluation systématique des noyaux d'inférence |
| EvoEngineer, arXiv:2510.03760 | évolution de code CUDA par LLM |
| ARGUS, arXiv:2604.18616 | optimisation agentique guidée par invariants de flux de données |
| Optimizing CUDA like a Human, arXiv:2606.26453 | micro-profileurs comme substituts d'expert |
| KForge, arXiv:2511.13274 | synthèse pour accélérateurs divers |
| Xe-Forge, arXiv:2605.26118 | optimisation multi-étages pour GPU Intel |
Le résultat à retenir : sur KernelBench, les modèles de raisonnement de pointe ne battent la ligne de base PyTorch que dans moins de 20 % des cas. Les approches agentiques avec profilage en boucle font nettement mieux.
3.7 Calcul scientifique¶
| Papier | Sujet |
|---|---|
| MFC 5.0: An exascale many-physics flow solver, arXiv:2503.07953 | CFD exascale sur GPU |
| Aurora: Architecting Argonne's First Exascale Supercomputer, arXiv:2509.08207 | architecture d'un supercalculateur exascale |
| GPU Performance Portability needs Autotuning, arXiv:2505.03780 | portabilité et réglage automatique |
| Mojo: MLIR-Based Performance-Portable HPC Science Kernels, arXiv:2509.21039 | Mojo en HPC |
3.8 Comment lire ces papiers¶
Une méthode en quatre passes
Passe 1 (5 min) — titre, résumé, figures, conclusion. Question : de quoi ça parle et quel est le résultat ?
Passe 2 (30 min) — introduction, méthode, résultats. Sauter les preuves et les détails d'implémentation. Question : comment font-ils ?
Passe 3 (2 h) — tout, en refaisant les calculs. Question : est-ce correct ?
Passe 4 — réimplémenter. Question : est-ce reproductible ?
Pour la plupart des papiers, la passe 1 suffit. Réservez les passes 3 et 4 aux papiers qui fondent votre travail.
Les chiffres des papiers ne sont pas des mesures indépendantes
Chaque accélération annoncée dépend de la ligne de base choisie, de sa version, du matériel, de la charge et de la configuration.
Vérifications utiles :
- quelle version de la ligne de base ? Un vLLM d'il y a un an n'est pas celui d'aujourd'hui ;
- quelle configuration ? Lot, longueur de contexte, précision ;
- la fourchette est-elle donnée ? « 1,0-1,7× » est plus informatif que « jusqu'à 1,7× » ;
- la comparaison au plancher physique est-elle faite ? MPK le fait (14,5 → 12,5 ms avec un plancher à ~10 ms), et c'est la meilleure pratique.
Chapitre suivant : 4 · Blogs et dépôts