spb/forge Public MIT

1<!-- Author: Simon-Pierre Boucher — contact@spboucher.ai -->23# Booster la vitesse d'inférence sur Apple Silicon — avec le format `.forge`4### Rapport de recherche web, août 2026 — llama.cpp/Metal, MLX, BaseRT, GGUF/safetensors/Core ML, Ollama/Xet56---78## TL;DR — la loi qui gouverne tout910**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 :11121. **Réduire les octets** → quantization 4-bit (~4× moins de trafic = ~4× plus de tok/s), KV f16, GQA132. **Saturer la bande passante** → GEMV avec déquantification fusionnée dans la boucle interne, lectures coalescées143. **Tuer l'overhead par token** → 1 command buffer/token, zéro allocation dans la boucle de décodage, kernels fusionnés1516Le *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).1718Ré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.1920---2122## 1. Le chemin de décodage rapide (kernels)2324- **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.25- **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).26- **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.27- **Routage par M** : la décision n°1 du dispatcher est mat-mat (prefill, M grand) vs mat-vec (décodage, M=1).2829## 2. KV cache3031- 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.32- **GQA est la plus grosse économie de bande passante KV** (déjà supporté par Forge via `n_kv_heads`).33- 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.3435## 3. Ce que la recherche valide dans `.forge` (et ce qu'elle ajoute)3637Notre 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.3839**Upgrades format à faire (issus de la recherche)** :40411. **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).422. **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).433. **MTLResidencySet** (macOS 15+) sur tous les buffers de poids, attaché à la queue — remplace `useResource`, évite les stalls d'éviction GPU.444. **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+).455. **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).466. **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`.4748## 4. À NE PAS faire (impasses documentées)4950- **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é.51- **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.52- **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.5354## 5. Roadmap Forge priorisée (impact ÷ effort)5556| # | Chantier | Gain attendu |57|---|---|---|58| 1 | **KV cache** (le `generate` actuel recompute tout le contexte à chaque token !) | 10–100× sur la génération, avant tout le reste |59| 2 | GEMV batch-1 f32→f16 + kernel matvec dédié (routage M=1) | ~2× décodage |60| 3 | Quant 4-bit planaire dans `.forge` + GEMV dequant-fusionné | ~4× décodage vs f32 |61| 4 | 1 command buffer/token, sampling GPU, zéro alloc | +15–20 % |62| 5 | Warmup séquentiel + ResidencySet + mlock opt-in au chargement `.forge` | cold-start ÷10, pas de stalls |63| 6 | Fusions (norm+matvec, SwiGLU, residual) | +15–50 % (résultat BaseRT) |64| 7 | Tenseurs ordre-forward + index repo + XXH3 footer dans `.forge` v2 | streaming + intégrité |65| 8 | M5 : `tensor_ops::matmul2d` pour le prefill (déjà exploré dans bench_precision) | ~4× prefill/TTFT |6667## Sources principales6869**Kernels/runtime** : [BaseRT](https://arxiv.org/html/2607.00501) · [BaseRT+M5](https://arxiv.org/html/2607.19438v1) · [llama.cpp Metal backend](https://deepwiki.com/ggml-org/llama.cpp/5.2-metal-backend-(apple)) · [MLX sur M5 (Apple ML)](https://machinelearning.apple.com/research/exploring-llms-mlx-m5) · [spec-decode Metal #23752](https://github.com/ggml-org/llama.cpp/issues/23752) · [MLX quantization](https://deepwiki.com/ml-explore/mlx/7-quantization) · [Rigel (simdgroup/fp8)](https://arxiv.org/abs/2606.12765) · [WWDC26 Metal tensors](https://developer.apple.com/videos/play/wwdc2026/330/)7071**Format/chargement** : [GGUF spec](https://github.com/ggml-org/ggml/blob/master/docs/gguf.md) · [safetensors](https://github.com/safetensors/safetensors) · [fastsafetensors](https://arxiv.org/html/2505.23072v1) · [repack au chargement PR #9921](https://github.com/ggml-org/llama.cpp/pull/9921) · [mmap (justine.lol)](https://justine.lol/mmap/) · [prefetch séquentiel #18758](https://github.com/ggml-org/llama.cpp/discussions/18758) · [Ollama blob store](https://deepwiki.com/ollama/ollama/2.4-storage-and-blob-transfer) · [HF Xet dedup](https://huggingface.co/docs/xet/en/deduplication) · [palettisation Core ML](https://apple.github.io/coremltools/docs-guides/source/opt-palettization-overview.html) · [BLAKE3/XXH3 vitesses](https://jolynch.github.io/posts/use_fast_data_algorithms/)72