La maggior parte dell’inference LLM in produzione passa da un percorso collaudato: GGUF, llama.cpp, fine. Funziona ed è veloce. Perché allora scrivere un runtime da zero? La risposta, spiega l’autore di un lungo tutorial su Towards Data Science, è l’ownership del decode stack: se devi cambiare il formato di quantization, aggiungere un sampler custom o una variante di attention, dentro una black box non puoi.
Il progetto, chiamato annotated-llm-runtime, è un motore CUDA/C++ per Qwen2.5-Coder-7B-Instruct su NVIDIA H100 (Hopper, sm_90), con ogni hot path commentato per spiegare il “perché” e non solo il “cosa”.
I numeri
Allo stato attuale il runtime raggiunge un decode steady-state di ~16,7 ms/token (~60 token/s) e un time-to-first-token di ~128 ms su un prompt da 512 token. Come riferimento, llama.cpp con Q4_K_M sulla stessa classe di GPU registra ~43 ms di TTFT e ~4,95 ms/token. Non è una gara: llama.cpp qui è il metro per capire se il proprio build si comporta in modo ragionevole.
Tre bug che insegnano
Il valore didattico sta in tre errori. Primo: un __syncthreads() dentro un branch warp-specialized è undefined behaviour — è una barriera a livello di blocco, e se solo un warp la raggiunge gli altri proseguono, corrompendo la softmax sulle page tail della KV cache. La soluzione è __syncwarp() dentro il branch e __syncthreads() fuori da qualsiasi controllo su warp-id.
Secondo: le CUDA graphs non sono un optional. Il decode eager, con centinaia di kernel launch per token, girava a ~119 ms/token; catturando la sequenza in un grafo e rieseguendola, il tempo è sceso a ~17 ms/token — un 7× a parità di kernel, solo eliminando l’overhead di submission verso il driver.
Terzo: sull’architettura Hopper la matematica dei formati conta. Per le GEMV di gate/up/down un unpack INT4 basato su prmt.b32 con attivazioni FP16 ha battuto la via __dp4a con attivazioni INT8, che è stata implementata, misurata e poi rimossa. “INT8 usa meno byte, quindi vince sul memory-bound” non è un teorema.
Il resto delle scelte è coerente con l’obiettivo di tenere piccolo lo spazio di debug: pesi INT4 simmetrici group-wise (gruppo 128) solo sulle sette proiezioni per layer, mentre embedding, rms_norm e lm_head restano in FP16. Il modello ha 28 layer, hidden size 3584 e una GQA con 28 query head e 4 KV head.
In sintesi
- Scrivere il proprio runtime serve a possedere il decode stack: quantization, sampler e attention diventano modificabili.
- Le CUDA graphs valgono un 7× (119 → 17 ms/token) senza toccare i kernel.
- Su Hopper le scelte di formato (INT4+FP16 vs INT8) vanno misurate, non date per scontate.
Fonte: How To Build Your Own LLM Runtime From Scratch — https://towardsdatascience.com/how-to-build-your-own-llm-runtime-from-scratch/
Hai qualcosa da aggiungere? Unisciti alla discussione.