4 · Tester et déboguer¶

Comment savoir qu'un noyau est correct, et quoi faire quand il ne l'est pas. Le GPU rend les deux plus difficiles que sur CPU.

4.1 Pourquoi c'est plus dur qu'ailleurs¶

Difficulté	Cause
Pas de `printf` utilisable	100 000 threads produiraient des mégaoctets
Pas de débogueur pas à pas pratique	`cuda-gdb` existe mais est lourd
Erreurs asynchrones	l'erreur remonte plusieurs appels plus tard
Erreurs collantes	une erreur corrompt tout le contexte
Non-déterminisme	l'ordre des blocs varie
Corruption silencieuse	une écriture hors borne ne plante pas toujours

La conséquence : sur GPU, la stratégie est la prévention et la vérification automatique, pas l'inspection.

4.2 Le test de correction¶

La structure minimale¶

import torch
import pytest

def reference(x, w):
    """Implémentation lente et évidemment correcte."""
    return torch.nn.functional.linear(x, w)

@pytest.mark.parametrize("m,n,k", [
    (1, 128, 128),          # cas minimal
    (17, 63, 129),          # dimensions non alignées
    (128, 128, 128),        # cas nominal
    (4096, 4096, 4096),     # grande taille
    (1, 1, 1),              # dégénéré
])
@pytest.mark.parametrize("dtype", [torch.float32, torch.bfloat16])
def test_mon_gemm(m, n, k, dtype):
    torch.manual_seed(0)
    x = torch.randn(m, k, device="cuda", dtype=dtype)
    w = torch.randn(n, k, device="cuda", dtype=dtype)

    obtenu = mon_gemm(x, w)
    attendu = reference(x, w)

    tol = {torch.float32: 1e-4, torch.bfloat16: 1e-2}[dtype]
    torch.testing.assert_close(obtenu, attendu, rtol=tol, atol=tol)

Les cinq cas à tester systématiquement

Le cas dégénéré : \(n = 1\), \(n = 0\).
Les dimensions non alignées : 17, 63, 129 — pas des puissances de 2. C'est là que les gardes de bord cassent.
Le cas nominal.
La grande taille : dépassements d'entiers 32 bits, quantification de vagues.
Les types : float32, bfloat16, float16 si supportés.

Le point 2 attrape la majorité des bugs réels.

Les tolérances¶

Précision de calcul	`rtol` raisonnable
`float64`	1e-10
`float32`	1e-5
TF32	1e-3
`bfloat16` / `float16`	1e-2
FP8	5e-2

Ces valeurs sont indicatives. La bonne tolérance dépend du nombre d'accumulations : l'erreur d'une somme de \(n\) termes croît en \(O(\sqrt{n}\,\varepsilon)\) pour des erreurs indépendantes.

Pour une GEMM \(K = 4096\) en BF16 accumulée en FP32 : \(\sqrt{4096}\times 1{,}2\times10^{-7} \approx 8\times10^{-6}\) — bien en dessous de la tolérance de \(10^{-2}\) imposée par la représentation finale.

Ne jamais utiliser assert torch.equal

Les réductions parallèles ne sont pas déterministes : l'ordre d'accumulation varie avec l'ordonnancement des blocs. Deux exécutions du même programme peuvent différer aux derniers bits.

Toujours torch.testing.assert_close avec une tolérance.

Les tests de propriétés¶

Complémentaires des tests de valeurs :

def test_linearite():
    x1, x2 = torch.randn(64, 128, device="cuda"), torch.randn(64, 128, device="cuda")
    w = torch.randn(256, 128, device="cuda")
    torch.testing.assert_close(mon_gemm(x1 + x2, w),
                               mon_gemm(x1, w) + mon_gemm(x2, w), rtol=1e-4)

def test_zero():
    x = torch.zeros(64, 128, device="cuda")
    w = torch.randn(256, 128, device="cuda")
    assert torch.all(mon_gemm(x, w) == 0)

def test_idempotence_relu():
    x = torch.randn(1000, device="cuda")
    torch.testing.assert_close(mon_relu(mon_relu(x)), mon_relu(x))

Ces tests attrapent des classes de bugs que les tests de valeurs manquent — notamment les erreurs d'indexation qui donnent un résultat plausible.

4.3 `compute-sanitizer`¶

L'outil le plus important de ce chapitre.

# Accès mémoire invalides
compute-sanitizer --tool memcheck ./prog

# Courses de données en mémoire partagée
compute-sanitizer --tool racecheck ./prog

# __syncthreads() dans des branches divergentes
compute-sanitizer --tool synccheck ./prog

# Mémoire non initialisée
compute-sanitizer --tool initcheck ./prog

Avec PyTorch :

compute-sanitizer --tool memcheck python mon_test.py

Sortie typique :

========= Invalid __global__ write of size 4 bytes
=========     at 0x70 in mon_noyau(float*, int)
=========     by thread (255,0,0) in block (39,0,0)
=========     Address 0x7f8e4a000400 is out of bounds

La règle

Tout noyau qui vous a coûté plus d'une heure de débogage doit passer par memcheck et racecheck.

Le ralentissement est de 10 à 100×, donc utilisez de petites tailles. Mais un accès hors borne trouvé en trente secondes économise deux jours.

4.4 La stratégie de bissection¶

Quand le résultat est faux et qu'on ne sait pas pourquoi.

Bissection sur les données¶

# Est-ce que ça marche avec des entrées triviales ?
x = torch.ones(...)          # tout à 1
x = torch.zeros(...)         # tout à 0
x = torch.eye(n)             # identité
x = torch.arange(n).float()  # valeurs identifiables

Avec arange, une erreur d'indexation devient visible : la sortie contient une valeur qui identifie l'entrée fautive.

Bissection sur la taille¶

for n in [1, 2, 31, 32, 33, 63, 64, 65, 127, 128, 129, 1024]:
    tester(n)

Les seuils 32, 64, 128 révèlent les bugs de warp et de tuile. Un noyau qui échoue à 33 mais pas à 32 a un problème de garde de bord.

Bissection sur le code¶

Neutraliser des parties du noyau :

__global__ void mon_noyau(...) {
    charger();
    __syncthreads();
    // calculer();          ← commenté : le chargement seul est-il correct ?
    ecrire_directement_ce_qui_a_ete_charge();
}

Chaque étape isolée, vérifiée, puis réintégrée.

Bissection sur la configuration¶

for threads in [32, 64, 128, 256]:
    for blocs in [1, 2, 7, 100]:
        tester(threads, blocs)

Un bug qui n'apparaît qu'avec plusieurs blocs est un problème de synchronisation inter-blocs ou d'atomiques.

4.5 Le journal sur GPU¶

Quand printf ne suffit plus.

struct Entree { int bloc, warp, etape; float valeur; };

__device__ Entree* g_journal;
__device__ unsigned  g_journal_idx;

__device__ void journaliser(int etape, float v) {
    if ((threadIdx.x & 31) == 0) {              // un thread par warp
        unsigned i = atomicAdd(&g_journal_idx, 1u);
        if (i < TAILLE_JOURNAL) {
            g_journal[i] = {(int)blockIdx.x, (int)(threadIdx.x >> 5), etape, v};
        }
    }
}

Côté hôte, on rapatrie et on analyse :

journal = journal_gpu.cpu().numpy()
df = pd.DataFrame(journal, columns=["bloc", "warp", "etape", "valeur"])
print(df.groupby("etape")["valeur"].describe())
# Chercher : NaN, valeurs aberrantes, étapes manquantes

C'est laborieux et c'est souvent le seul moyen de comprendre ce qui se passe dans un noyau complexe — un megakernel notamment.

4.6 Les symptômes et leurs causes¶

Symptôme	Causes probables, par ordre
NaN dans la sortie	division par zéro, `exp` qui déborde, mémoire non initialisée, race
Faux uniquement en grande taille	dépassement d'`int32` (`n*m` > 2³¹), garde de bord
Faux uniquement à plusieurs blocs	synchronisation inter-blocs, atomique manquant
Faux uniquement en `-O3`	variable non `volatile`, `__threadfence()` manquant
Non déterministe	course de données, atomique manquant, ordre de réduction
Correct sur GPU A, faux sur GPU B	taille de warp (AMD), mémoire partagée disponible, TF32
Lent sans raison	spilling, contention atomique, accès non coalescés
`illegal memory access`	hors borne — `memcheck` le localise
Le noyau ne se termine pas	boucle d'attente, non-résidence des blocs

Le piège du dépassement d'entier¶

// BUG : n et m sont des int, leur produit déborde à 2^31
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float v = donnees[idx * m + j];        // ← déborde si idx * m > 2 milliards

// CORRECT
size_t idx = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
float v = donnees[idx * m + j];

Un tenseur de \(2^{31}\) éléments fait 8 Go en float32 — parfaitement courant. Ce bug est fréquent et se manifeste par des accès à des adresses aberrantes.

4.7 Le test de non-régression de performance¶

La correction ne suffit pas : une optimisation peut casser silencieusement les performances.

import json, pathlib

REFERENCE = pathlib.Path("bench_reference.json")

def test_performance():
    x = torch.randn(1 << 24, device="cuda")
    t = mesurer(mon_noyau, x)

    if REFERENCE.exists():
        ref = json.loads(REFERENCE.read_text())["mon_noyau"]
        assert t < ref * 1.10, (
            f"Régression : {t*1000:.1f} µs contre {ref*1000:.1f} µs de référence")
    else:
        REFERENCE.write_text(json.dumps({"mon_noyau": t}))

La marge de 10 % n'est pas arbitraire

Sans verrouillage des fréquences, la variabilité entre deux exécutions atteint ±15 %. Un seuil trop serré produit des faux positifs permanents.

Sur une machine de test dédiée avec fréquences verrouillées, on peut descendre à 5 %.

4.8 La checklist de débogage¶

Quand un noyau donne un résultat faux :

□  1. compute-sanitizer --tool memcheck          → accès hors borne ?
□  2. compute-sanitizer --tool racecheck         → course en mémoire partagée ?
□  3. Tester avec n = 1, puis 32, 33, 64, 65     → garde de bord ?
□  4. Tester avec 1 seul bloc                    → problème inter-blocs ?
□  5. Entrée = arange                            → erreur d'indexation ?
□  6. Neutraliser le calcul, vérifier le chargement seul
□  7. Vérifier les types (int vs size_t) et les dépassements
□  8. Vérifier les tolérances du test elles-mêmes
□  9. Comparer -G et -O3                         → volatile manquant ?
□ 10. Journal sur GPU en dernier recours

Résumé du chapitre¶

À retenir

Sur GPU, la stratégie est la prévention et la vérification automatique, pas l'inspection.
Toujours torch.testing.assert_close avec tolérance, jamais torch.equal : les réductions ne sont pas déterministes.
Tester cinq cas : dégénéré, dimensions non alignées (17, 63, 129), nominal, grande taille, plusieurs types. Le second attrape la majorité des bugs.
Les tests de propriétés (linéarité, zéro, idempotence) attrapent ce que les tests de valeurs manquent.
compute-sanitizer avec memcheck et racecheck est l'outil le plus rentable du chapitre.
Bissection sur les données, la taille, le code et la configuration.
Attention aux dépassements d'int32 : un tenseur de 8 Go les déclenche.
Test de non-régression de performance avec 10 % de marge (5 % à fréquences verrouillées).

Vérifiez que vous avez compris¶

Votre noyau donne des NaN uniquement avec des entrées de grande magnitude. Où chercher ?

Presque certainement un débordement d'exponentielle.

En float32, expf(x) déborde pour \(x > 88\). En float16, dès \(x > 11\). Un débordement donne inf, puis inf - inf ou inf / inf donne NaN.

Le remède standard est la soustraction du maximum avant l'exponentielle :

float m = reduce_max(x);
float e = __expf(x - m);      // toujours ≤ 1, jamais de débordement

C'est exactement ce que fait le softmax stable, et par extension FlashAttention. Voir IA 3.

Autres causes possibles : division par une somme nulle, sqrt d'un nombre négatif dû à une erreur d'arrondi dans une variance.

Votre noyau passe tous les tests mais casse en production. Quel test manquait ?

Le plus probable : les dimensions réelles.

Les tests utilisent souvent des puissances de deux (128, 256, 1024). La production utilise des dimensions arbitraires : une taille de vocabulaire de 128 256, une dimension cachée de 2 880, un lot de 37.

Les bugs typiques que cela révèle :

gardes de bord incorrectes sur la dernière tuile ;
hypothèses d'alignement violées (float4 sur une adresse non multiple de 16) ;
contraintes d'un noyau quantifié — l'issue vLLM #38022 documente exactement cela : le noyau Marlin MoE échoue avec \(K = N = 2880\) parce qu'il exige des multiples de 128.

Le remède : inclure les dimensions réelles de votre modèle dans la paramétrisation des tests, et pas seulement des puissances de deux.

racecheck signale une course mais le résultat semble toujours correct. Faut-il corriger ?

Oui, absolument.

Une course de données est un comportement indéfini. Qu'elle ne se manifeste pas aujourd'hui dépend de facteurs qui changeront :

l'ordonnancement des warps (varie avec la charge du GPU) ;
la version du compilateur (le réordonnancement d'instructions change) ;
l'architecture (les latences relatives diffèrent) ;
la taille du problème (le nombre de warps résidents change).

Un bug qui se manifeste une fois sur \(10^6\) sur votre machine peut se manifester une fois sur \(10^3\) en production, ou systématiquement sur la génération suivante de matériel.

C'est le type de bug le plus coûteux : il apparaît tard, sous charge, et il est très difficile à reproduire. Corrigez-le pendant que racecheck vous dit exactement où il est.

Chapitre suivant : 5 · Plan d'entraînement de 90 jours