// data & ai · giornale tecnico MILANO ● LIVE 00:00:00

Costruire un runtime di inference LLM da zero: le lezioni di un motore CUDA per Qwen2.5-Coder-7B

Un motore CUDA da zero per Qwen2.5-Coder-7B su H100: perché possedere il decode stack, i numeri rispetto a llama.cpp e i tre bug (syncthreads, CUDA graphs, INT4 vs INT8) che insegnano di più.

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/

Condividi X Facebook LinkedIn WhatsApp Email

// scritto da

Fernando

Hai qualcosa da aggiungere? Unisciti alla discussione.

Il tuo indirizzo email non sarà pubblicato. I campi obbligatori sono contrassegnati *

Altri dell'autore

dalla stessa firma