Aller au contenu

5 · La quantification

Diviser par deux ou par quatre le volume des poids, sur un problème limité par la bande passante. C'est le levier le plus direct de l'inférence — et celui qui demande le plus de soin.

5.1 Pourquoi ça marche si bien en inférence

Rappel du chapitre 1 : en décodage à lot 1, le temps est dicté par la lecture des poids.

\[ t_{\min} = \frac{\text{octets des poids}}{B} \]

Diviser les octets par 4 divise le temps par ~4. Aucune optimisation de calcul ne peut faire cela.

Exemple : Llama-70B sur H100 (3,35 To/s), en décodage à lot 1.

Précision Poids \(t_{\min}\)/jeton Jetons/s max
BF16 140 Go 41,8 ms 24
FP8 70 Go 20,9 ms 48
FP4 / INT4 35 Go 10,4 ms 96

(140 Go ne tenant pas sur un H100 80 Go, la ligne BF16 suppose deux cartes ; les ordres de grandeur restent valides.)

Second effet : sur les GPU récents, le débit tensor core double aussi à chaque division par deux de la largeur (voir Fondations 5). En préremplissage, où l'on est limité par le calcul, on gagne donc également.

5.2 Les schémas de quantification

La notation standard est W\(x\)A\(y\) : poids sur \(x\) bits, activations sur \(y\) bits.

Schéma Poids Activations Usage
W16A16 BF16 BF16 référence
W8A8 FP8 FP8 inférence, matériel Hopper+
W8A16 FP8 ou INT8 BF16 GPU sans FP8 natif (Ampere)
W4A16 INT4 / FP4 BF16 décodage à petit lot — le plus répandu
W4A8 4 bits FP8 agressif
NVFP4 / MXFP4 FP4 + échelle par bloc FP8 ou BF16 Blackwell, CDNA4

Pourquoi W4A16 domine en décodage local

Parce que le goulot est la lecture des poids, pas celle des activations.

À lot 1, les activations font quelques kilooctets ; les poids font des gigaoctets. Quantifier les activations n'apporte presque rien et dégrade la qualité. Quantifier les poids apporte tout.

À grand lot, l'arbitrage change : les activations deviennent volumineuses et le calcul devient limitant, ce qui favorise W8A8.

5.3 Ce qui se passe dans le noyau

Un noyau W4A16 fait quelque chose de contre-intuitif : il déquantifie les poids avant de les multiplier.

Poids INT4 en HBM
    │  lecture : 4 bits par poids ← LE GAIN EST ICI
    ▼
Déquantification en registres : w = (q - z) * s
    │
    ▼
Poids BF16 en registres
    │
    ▼
Tensor core (BF16 × BF16 → FP32)

Le calcul se fait donc en 16 bits. Le gain est entièrement dans le transfert.

Le coût : la déquantification consomme des instructions arithmétiques. Sur un noyau limité par la mémoire, elles sont gratuites (elles se cachent derrière les chargements) ; c'est la clé de la viabilité du schéma.

C'est précisément ce que fait Marlin, le noyau de référence : il « masque le surcoût de déquantification en recouvrant efficacement la latence d'accès aux données avec les opérations flottantes de la déquantification et du calcul de la GEMM ».

Sur Blackwell et CDNA 4, les instructions tcgen05.mma et MFMA acceptent directement des opérandes FP4/FP6 avec échelle par blocs : la déquantification explicite disparaît, câblée dans le tensor core.

5.4 Les facteurs d'échelle

Un entier 4 bits représente 16 valeurs. Pour couvrir une plage utile, il faut un facteur d'échelle.

\[ w \approx s \cdot (q - z) \]

où \(q\) est l'entier quantifié, \(s\) l'échelle (flottant), \(z\) le point zéro.

La granularité de \(s\) décide de la qualité :

Granularité \(s\) par… Surcoût mémoire Qualité
Par tenseur matrice entière ~0 mauvaise
Par canal ligne ou colonne ~0,1 % bonne
Par groupe de 128 bloc de 128 poids ~1,5 % très bonne
Par bloc de 32 (MX) bloc de 32 poids ~3 % excellente
Par bloc de 16 (NVFP4) bloc de 16 poids ~6 % excellente

Les standards MX (microscaling) de l'Open Compute Project définissent MXFP8, MXFP6 et MXFP4 avec une échelle E8M0 (un exposant sur 8 bits) par bloc de 32. NVFP4, la variante NVIDIA, utilise une échelle FP8 E4M3 par bloc de 16.

Les deux sont supportés en matériel par Blackwell ; CDNA 4 supporte les formats MX.

5.5 Les méthodes de quantification

Comment choisir les \(q\) et les \(s\) ?

Arrondi au plus proche (RTN)

Le plus simple : \(q = \text{round}(w/s)\). Rapide, sans calibration, et suffisant en 8 bits. Dégrade sensiblement en 4 bits.

GPTQ

Quantifie couche par couche, colonne par colonne, en compensant l'erreur introduite sur les colonnes restantes via une approximation de la hessienne. Nécessite un petit jeu de calibration.

AWQ (Activation-aware Weight Quantization)

Observation : une petite fraction des poids (~1 %) contribue de façon disproportionnée à la sortie. AWQ identifie les canaux importants à partir de la magnitude des activations (pas des poids) et les protège par une mise à l'échelle par canal.

SmoothQuant

Déplace la difficulté des activations vers les poids par une transformation mathématiquement équivalente :

\[\mathbf{X}\mathbf{W} = (\mathbf{X}\operatorname{diag}(\mathbf{s})^{-1})(\operatorname{diag}(\mathbf{s})\mathbf{W})\]

Utile pour W8A8, où les valeurs aberrantes d'activation sont le problème.

QAT (Quantization-Aware Training)

Simuler la quantification pendant l'entraînement ou l'affinage. Meilleure qualité, coût d'entraînement.

Ce qu'il faut retenir en pratique

  • 8 bits : RTN suffit presque toujours.
  • 4 bits : GPTQ ou AWQ, avec une calibration sur quelques centaines d'échantillons représentatifs.
  • En dessous de 4 bits : QAT devient nécessaire, et la dégradation redevient significative.

5.6 Les noyaux de référence

Noyau Schéma Particularité
Marlin W4A16 recouvre la déquantification avec la GEMM ; référence sur Ampere/Ada
Machete W4A16 successeur de Marlin pour Hopper
Marlin MoE W4A16 groupé pour les MoE
DeepGEMM (DeepSeek) FP8 avec échelle fine GEMM et GEMM groupée
bitsandbytes NF4, INT8 utilisé par QLoRA
TensorRT-LLM tous intégré au moteur
Transformer Engine FP8 gestion automatique des échelles à l'entraînement

Les contraintes d'alignement des noyaux quantifiés

Ces noyaux imposent des contraintes de dimensions. Un bogue rapporté en 2026 sur vLLM illustre le problème : le noyau Marlin MoE échoue sur un modèle MXFP4 avec \(K = N = 2880\), parce qu'il exige des dimensions alignées sur 128.

C'est un piège courant : un modèle dont les dimensions ne sont pas des multiples de 128 (ou 64, selon le noyau) peut ne pas être servable avec le noyau quantifié le plus rapide, et retomber sur un chemin lent — ou planter.

Vérifiez les dimensions de votre modèle avant de choisir un schéma.

5.7 Le cache KV

Souvent oublié, et il représente une part croissante du trafic à contexte long.

Quantifier le cache en FP8 ou INT8 :

  • divise par 2 son empreinte et le trafic associé ;
  • permet des contextes deux fois plus longs à mémoire égale ;
  • avec une dégradation généralement faible, mais à vérifier — le cache est lu à chaque jeton, donc l'erreur se propage.

vLLM et SGLang supportent kv_cache_dtype="fp8". Sur des contextes très longs (> 100 k jetons), c'est souvent le levier le plus rentable, puisque le cache dépasse alors les poids.

5.8 Ce que la quantification casse

Les six pièges

1. Les couches sensibles. Plongements, dernière couche (lm_head), normalisations. On les garde en 16 bits. Coût mémoire faible, gain de qualité important.

2. Les valeurs aberrantes d'activation. Certains canaux ont des activations 100× plus grandes que les autres. En W8A8, elles saturent. SmoothQuant et AWQ existent pour ça.

3. La calibration non représentative. GPTQ et AWQ calibrent sur un jeu d'échantillons. Si ce jeu ne ressemble pas à votre usage réel (autre langue, autre domaine, autre longueur), la quantification est mal réglée.

4. L'évaluation trompeuse. La perplexité bouge peu ; les capacités de raisonnement, de suivi d'instructions et de code peuvent chuter nettement. Évaluez sur ce qui vous importe, pas sur la perplexité.

5. Le gain qui n'arrive pas. Si votre noyau quantifié n'est pas optimisé, la déquantification peut coûter plus que le transfert économisé. Mesurez.

6. L'entraînement. La quantification agressive est une technique d'inférence. L'entraînement en 4 bits reste un sujet de recherche ; l'entraînement en FP8 est en revanche déployé (Transformer Engine, DeepSeek).

Résumé du chapitre

À retenir

  • En décodage, le temps est dicté par la lecture des poids. Diviser leur volume par 4 divise le temps par ~4.
  • Notation W\(x\)A\(y\). W4A16 domine le décodage à petit lot, W8A8 le service à grand lot.
  • Un noyau W4A16 déquantifie en registres avant la MMA : le gain est entièrement dans le transfert, et la déquantification se cache derrière les chargements (c'est le principe de Marlin).
  • Sur Blackwell et CDNA 4, tcgen05/MFMA acceptent directement des opérandes FP4 avec échelle par blocs.
  • La granularité de l'échelle décide de la qualité : par bloc de 32 (MX) ou 16 (NVFP4).
  • 8 bits : RTN suffit. 4 bits : GPTQ ou AWQ avec calibration.
  • Quantifier le cache KV est souvent le levier le plus rentable au-delà de 100 k jetons de contexte.
  • Vérifiez les contraintes d'alignement des noyaux (Marlin exige des multiples de 128).

Vérifiez que vous avez compris

Pourquoi W4A16 plutôt que W4A4 en décodage à lot 1 ?

Parce que les activations ne pèsent rien à lot 1.

Pour un modèle de dimension \(d = 8192\), l'activation d'une couche fait \(8192 \times 2 = 16\) Ko en BF16. Les poids d'une couche font des centaines de mégaoctets. Quantifier les activations économise 8 Ko sur des centaines de mégaoctets : négligeable.

En revanche, W4A4 dégraderait sensiblement la qualité (les activations ont des valeurs aberrantes) et exigerait des noyaux plus complexes.

À grand lot, le calcul change : les activations deviennent \(b \times 16\) Ko et le régime devient limité par le calcul, ce qui rend W8A8 (voire W4A8) intéressant.

Votre modèle W4A16 est plus lent que le W16A16. Comment est-ce possible ?

Plusieurs causes possibles, toutes réelles :

  1. Le noyau n'est pas optimisé. Une déquantification naïve, non recouverte avec les chargements, ajoute des instructions sans rien économiser. Vérifiez que vous utilisez Marlin ou Machete, et non un chemin de repli.
  2. Les dimensions ne conviennent pas. Si \(K\) ou \(N\) n'est pas aligné comme le noyau l'exige, le moteur retombe sur une implémentation lente.
  3. Vous êtes limité par le calcul, pas par la mémoire. En préremplissage ou à grand lot, la quantification des poids n'aide pas si le calcul se fait toujours en 16 bits.
  4. Le cache KV domine. À contexte long, quantifier les poids ne change presque rien si 80 % du trafic est le cache.

Diagnostic : dram__bytes.sum avant et après. S'il n'a pas baissé, c'est la cause 4 ou 3.

Pourquoi l'échelle par bloc de 32 est-elle si supérieure à l'échelle par tenseur ?

Parce que la dynamique des poids varie énormément d'une région à l'autre de la matrice.

Avec une échelle unique, elle doit couvrir le poids de plus grande magnitude de toute la matrice. Tous les poids d'une région à faible dynamique sont alors quantifiés sur une fraction des 16 niveaux disponibles — parfois 2 ou 3 niveaux effectifs.

Avec une échelle par bloc de 32, chaque bloc utilise la pleine étendue des 16 niveaux pour sa propre dynamique locale.

Le coût est modeste : une échelle E8M0 (8 bits) pour 32 poids de 4 bits, soit \(8/(32 \times 4) = 6{,}25\ \%\) de surcoût — en pratique ~3 % après optimisations d'encodage. C'est ce compromis qui rend le 4 bits utilisable, et c'est pourquoi il a été normalisé (MX) puis câblé en matériel.

Chapitre suivant : 6 · MoE et GEMM groupée

Sources de ce chapitre