# Booster la vitesse d'inférence sur Apple Silicon — avec le format .forge

# Rapport de recherche web, août 2026 — llama.cpp/Metal, MLX, BaseRT, GGUF/safetensors/Core ML, Ollama/Xet

# TL;DR — la loi qui gouverne tout

Le décodage autorégressif (batch 1) est borné par la bande passante mémoire, pas par le calcul : latence/token ≈ (octets de poids + KV lus par token) ÷ bande passante. Budgets M-series : M4 120 GB/s · M5 153 · M4 Pro 273 · M5 Max ≈ 550 · M2/M3 Ultra ~800. Trois leviers, dans l'ordre :

1. Réduire les octets → quantization 4-bit (~4× moins de trafic = ~4× plus de tok/s), KV f16, GQA
2. Saturer la bande passante → GEMV avec déquantification fusionnée dans la boucle interne, lectures coalescées
3. Tuer l'overhead par token → 1 command buffer/token, zéro allocation dans la boucle de décodage, kernels fusionnés

Le prefill, lui, est borné par le calcul → simdgroup_matrix (déjà dans Forge) et, sur M5, les Neural Accelerators via tensor_ops::matmul2d (~4× le prefill de M4 dans MLX ; le décodage ne gagne que 19–27 % — bande passante toujours).

Références de vitesse : Llama-3.1-8B Q4 ≈ 51–75 tok/s sur M4 Max ; BaseRT (meilleur runtime Metal publié) : Qwen3-0.6B Q4 à 464 tok/s sur M4 Pro, +15–56 % vs llama.cpp, principalement par fusion de kernels.

# 1. Le chemin de décodage rapide (kernels)

• GEMV dequant-fusionné : le kernel matvec déquantifie les blocs 4-bit dans la boucle interne (jamais de poids f32 matérialisés). Design llama.cpp : chaque simdgroup possède une bande de lignes de sortie (N_R0 lignes/simdgroup, N_SG simdgroups/threadgroup, réglés par format) ; 32 lanes stride le long des blocs de la ligne ; réduction simd_sum finale. Intensité arithmétique ~2 ops/octet — pur streaming.
• Fusions à implémenter (chaque fusion = un lancement + un aller-retour mémoire en moins ; l'overhead de dispatch coûtait +17 % sur SmolLM2-360M avant fusion) : rmsnorm+matvec QKV, residual+norm, gate·up SwiGLU en un kernel, dequant+matvec (toujours), matvec logits + sampling top-k sur GPU (ne rapatrier que le token).
• Zéro overhead par token : un seul command buffer encodant toute la pile de couches ; jamais de waitUntilCompleted au milieu ; double-buffering (encoder le token N+1 pendant que N s'exécute) ; « the decode loop allocates zero bytes » (BaseRT). Spécialiser par MTLFunctionConstantValues (head_dim, n_heads, params quant cuits dans le binaire) — Forge fait déjà ça pour le matmul.
• Routage par M : la décision n°1 du dispatcher est mat-mat (prefill, M grand) vs mat-vec (décodage, M=1).

# 2. KV cache

• Pré-alloué à la longueur max au chargement, layout coalescé le long de l'axe temporel par tête ([layer][head][seq][d]), f16 par défaut. Kernel décodage dédié : une query, split-K le long du cache pour les longs contextes.
• GQA est la plus grosse économie de bande passante KV (déjà supporté par Forge via n_kv_heads).
• Paged attention sur Metal : existe en expérimental, personne ne le met en prod pour du batch-1 — cache contigu simple = plus rapide. KV quantifié q8/q4 : fragile sur Metal (bugs llama.cpp), pas prioritaire.

# 3. Ce que la recherche valide dans .forge (et ce qu'elle ajoute)

Notre choix central — tenseurs alignés page 16 KB + mmap + newBuffer(bytesNoCopy) — est exactement le pattern llama.cpp/Metal, en mieux : GGUF n'aligne qu'à 32 octets (obligeant des fenêtres page-alignées bricolées) et safetensors à 8 octets — c'est l'anti-pattern démontré (impossible de passer un tenseur à bytesNoCopy). Le sharding contenu-adressé + manifests ≈ le blob store d'Ollama et les xorbs de HF Xet — l'industrie a convergé sur notre design.

Upgrades format à faire (issus de la recherche) :

1. Ordonner les tenseurs dans l'ordre du forward pass dans les shards (embeddings → blk.0 … blk.N → head) : le streaming de couches devient de l'I/O séquentielle (prefetch du shard N+1 pendant le calcul de la couche N, ~+10 % façon AirLLM).
2. Warmup séquentiel au chargement : le demand-faulting à froid est lent (fautes 16 KB aléatoires) ; un passage F_RDADVISE/lecture séquentielle par shard sature le SSD (~7 GB/s) — llama.cpp mesure +13–14 %. Option mlock quand le modèle tient en RAM (l'éviction du page cache sous pression = 0.025 tok/s de thrash).
3. MTLResidencySet (macOS 15+) sur tous les buffers de poids, attaché à la queue — remplace useResource, évite les stalls d'éviction GPU.
4. Quantization planaire, pas interleavée : data 4-bit packée en uint32 + tenseurs scales/biases séparés (style MLX, groupe 64) plutôt que les blocs 18 octets de GGUF — garde les loads vectoriels 16 B alignés et les scales résidents en threadgroup memory. Ajouter un type LUT/palettisé (indices 4-bit + LUT f16 16 entrées, style Core ML). Réserver des IDs dtype fp8 (E4M3/E5M2 — footprint seulement : fp8 émulé à 0.94× f16 sur M4, natif sur M5+).
5. Layout canonique sur disque : la seule transformation on-disk qui paie est la pré-transposition pour le sens d'accès du matvec ; PAS de fichiers pré-swizzlés (llama.cpp a supprimé Q4_0_4_4 et repacke au chargement — mais sur Apple Silicon, préférer des kernels qui consomment le layout canonique, car repacker touche toutes les pages et tue le zéro-copie).
6. Intégrité en deux vitesses : BLAKE3 (7–16 GB/s, Merkle) comme adresse de contenu vérifiée au pull/push ; XXH3 (30 GB/s) par tenseur en footer pour un --verify opt-in fusionné au warmup (coût ~0). Jamais de hash sur le chemin de chargement par défaut — mmap veut dire qu'on n'a pas encore lu les octets. + GC mark-and-sweep manifests→shards (style Ollama) et un index repo tensor → (shard, offset) façon safetensors.index.json.

# 4. À NE PAS faire (impasses documentées)

• Kernels ternaires GPU : TQ1_0/TQ2_0 ciblent les CPU ; pas d'implémentation Metal, jugée difficile par les mainteneurs. Notre QAT ternaire reste un outil d'entraînement/recherche ; à l'inférence, servir en 4-bit groupé.
• Speculative decoding naïf : perte nette sur Metal à toutes les configs (le coût de dispatch draft/verify mange le gain). Ne marche que verify-fusionné dans un seul graphe (Recurrent Drafter d'Apple : 2.3×) — plus tard.
• ANE pour le décodage : Core ML impose 2–4× d'overhead, shapes statiques ; le GPU est 2–5× plus rapide. Niche : basse conso, ou prefill-ANE + decode-GPU.

# 5. Roadmap Forge priorisée (impact ÷ effort)

# Sources principales

Kernels/runtime : BaseRT · BaseRT+M5 · llama.cpp Metal backend · MLX sur M5 (Apple ML) · spec-decode Metal #23752 · MLX quantization · Rigel (simdgroup/fp8) · WWDC26 Metal tensors

Format/chargement : GGUF spec · safetensors · fastsafetensors · repack au chargement PR #9921 · mmap (justine.lol) · prefetch séquentiel #18758 · Ollama blob store · HF Xet dedup · palettisation Core ML · BLAKE3/XXH3 vitesses