Contents
- Step 00Orientation
- Step 01The front door
- Step 02Configuration
- Step 03The engine loop: LLMEngine
- Step 04The unit of work: Sequence
- Step 05Paged KV memory: BlockManager
- 5.1Blocks, block tables and slots
- 5.2Block and the four bookkeeping structures
- 5.3Chained block hashes
- 5.4can_allocate: count cache hits, check capacity
- 5.5allocate: build the block table
- 5.6_allocate_block: take a block, evict its old identity
- 5.7deallocate: release a sequence's blocks
- 5.8can_append and may_append: growing during decode
- 5.9hash_blocks: publish newly full blocks
- Step 06Choosing each batch: Scheduler
- Step 07From sequences to tensors: ModelRunner
- 7.1Boot sequence
- 7.2The control plane: call, write_shm, read_shm, loop
- 7.3warmup_model: measure the worst prefill
- 7.4allocate_kv_cache: size the pool from what is left
- 7.5prepare_prefill: pack a varlen batch
- 7.6prepare_decode: one token per sequence
- 7.7prepare_block_tables and prepare_sample
- 7.8run and run_model: eager or graph replay
- 7.9capture_cudagraph: one graph per batch bucket
- 7.10exit
- Step 08The side channel: Context
- Step 09The model: models/qwen3.py
- Step 10The layers: nanovllm/layers/
- 10.1linear.py: tensor-parallel linear layers
- 10.2embed_head.py: vocab-parallel embedding and LM head
- 10.3attention.py: storing K/V and calling FlashAttention
- 10.4layernorm.py: RMSNorm with a fused residual add
- 10.5rotary_embedding.py: RoPE
- 10.6activation.py: SiluAndMul
- 10.7sampler.py: temperature and the exponential race
- Step 11Weight loading: utils/loader.py
- Step 12One step, end to end, with real shapes
- Step 13Where each optimisation lives
- Step 14Sharp edges and limits
- Step 15Glossary and next exercises
Step 00Orientation
Before reading any function, get three things straight: which files exist, how they stack, and the five ideas from LLM serving that every one of them assumes.
0.1Repository map
Twenty-one Python files. The package is nanovllm/; example.py and bench.py are the two programs that use it. Every file name below links to the step that explains it. Line counts include blank lines.
| File | Lines | Role |
|---|---|---|
| README.md | 66 | Install, model download, quick start, benchmark table |
| pyproject.toml | 27 | Package metadata and the five runtime dependencies |
| example.py | 33 | Minimal chat-style usage |
| bench.py | 32 | Throughput benchmark: 256 random requests |
| nanovllm/__init__.py | 2 | Exports LLM and SamplingParams |
| nanovllm/llm.py | 5 | LLM is an empty subclass of LLMEngine |
| nanovllm/config.py | 25 | Every engine knob, validated in __post_init__ |
| nanovllm/sampling_params.py | 11 | temperature, max_tokens, ignore_eos |
| nanovllm/engine/llm_engine.py | 90 | Builds everything; generate() → step() loop |
| nanovllm/engine/sequence.py | 83 | One request: tokens, counters, block table |
| nanovllm/engine/block_manager.py | 120 | Paged KV allocator with prefix caching |
| nanovllm/engine/scheduler.py | 92 | Chooses each step's batch; preemption; bookkeeping |
| nanovllm/engine/model_runner.py | 257 | Per-GPU worker: KV pool, tensors, CUDA graphs, IPC |
| nanovllm/utils/context.py | 27 | Global per-step metadata read inside the model |
| nanovllm/utils/loader.py | 28 | safetensors → sharded parameters |
| nanovllm/models/qwen3.py | 216 | Qwen3 architecture assembled from layers |
| nanovllm/layers/linear.py | 156 | Replicated / column / row / merged / QKV parallel linears |
| nanovllm/layers/embed_head.py | 66 | Vocab-parallel embedding and LM head |
| nanovllm/layers/attention.py | 75 | Triton KV-store kernel + FlashAttention calls |
| nanovllm/layers/layernorm.py | 50 | RMSNorm with fused residual add |
| nanovllm/layers/rotary_embedding.py | 59 | RoPE cos/sin cache |
| nanovllm/layers/activation.py | 11 | SiluAndMul |
| nanovllm/layers/sampler.py | 12 | Temperature + exponential-race sampling |
The .gitignore is the stock GitHub Python template, LICENSE is MIT, and assets/logo.png is the README logo. None of them affect behaviour.
0.2How the pieces stack
Context is a module-level global that the runner fills before each forward and that Attention and ParallelLMHead read.Read the picture as a call stack. Your script calls LLM.generate. The engine decides, on the CPU, which sequences run in the next step and where their keys and values live in GPU memory. The runner turns that decision into tensors, runs the model, and returns one sampled token per sequence. The model is an ordinary PyTorch module tree whose leaves are the layer files.
0.3Five ideas the code assumes
Each idea below names where it lives in this repository, so you can jump ahead if you already know it.
- Prefill and decode
- A request starts with a prefill: the whole prompt goes through the model in one forward pass, which writes every prompt token's keys and values (K/V) into a cache and produces the first output token. After that, each decode step feeds exactly one new token per sequence. Prefill is compute-bound and has many tokens per sequence; decode is memory-bound and has one. nano-vllm never mixes the two inside one step (Step 06).
- KV cache
- Attention for token t needs the K/V of every earlier token, at every layer. Keeping them on the GPU means a decode step only computes K/V for the one new token. For Qwen3-0.6B that cache costs 112 KiB per token (Step 07.4).
- Paged KV
- The cache is cut into fixed-size blocks of 256 token slots. A sequence owns a list of block ids, its block table, which do not have to be adjacent. That removes fragmentation and lets two sequences point at the same block when they share a prompt prefix (Step 05).
- Continuous batching
- The batch is rebuilt every step. Finished sequences leave and waiting ones join between steps, so the GPU stays busy without waiting for the slowest request in a fixed batch (Step 06).
- Tensor parallelism
- With
tensor_parallel_size = N, N processes each hold 1/N of most weight matrices and 1/N of the attention heads, and combine partial results with NCCL collectives (Step 10.1).
0.4The route through this guide
- Steps 01–03 cover the surface: how the package is used, every configuration knob, and the engine loop that drives everything.
- Steps 04–06 are the CPU-side bookkeeping: the
Sequenceobject, the paged block allocator with prefix caching, and the scheduler. Step 06 ends with a simulator that runs a JavaScript port of those three files. - Steps 07–08 cross to the GPU: the per-rank
ModelRunnerthat sizes the KV cache, builds the batch tensors and replays CUDA graphs, and theContextside channel it fills. - Steps 09–11 are the model: Qwen3's module tree, each layer file, and the weight loader that fills sharded parameters.
- Steps 12–15 put it back together: one step traced with real shapes, an index of where each optimisation lives, sharp edges found while reading, and a glossary.
Code blocks show the exact lines from commit e86692c with their real line numbers; highlighted lines are the ones the text discusses. Use Copy for the code without line numbers, or GitHub ↗ to open the same range in the repository. The “Mark this step as read” buttons tick the route in the sidebar; that progress is kept only in this browser.
Step 01The front door
What a user of the package touches: two exported names, two example programs, and the packaging metadata.
1.1The public surface: LLM and SamplingParams
The package exports exactly two names. LLM is an empty subclass of LLMEngine; it exists so that from nanovllm import LLM reads the same as vLLM's from vllm import LLM. All behaviour is in LLMEngine (Step 03).
1from nanovllm.llm import LLM2from nanovllm.sampling_params import SamplingParams1from nanovllm.engine.llm_engine import LLMEngine2 3 4class LLM(LLMEngine):5 pass1.2example.py: the smallest real use
- The model path must be a local directory.
Configassertsos.path.isdir(model), so a Hugging Face hub id such as"Qwen/Qwen3-0.6B"will not work. The README'shuggingface-cli download … --local-dir ~/huggingface/Qwen3-0.6B/command produces the directory this script expects. - The script loads its own tokenizer only to apply Qwen3's chat template with
tokenize=False, so the prompts it passes are strings. The engine tokenizes them again insideadd_request. enforce_eager=Trueskips CUDA graph capture. Start-up is faster; decode steps are slower (Step 07.9).- One
SamplingParamsobject is shared by both prompts;generatebroadcasts it. - Each output is a dict with
"text"and"token_ids", in the same order as the prompts.
1import os2from nanovllm import LLM, SamplingParams3from transformers import AutoTokenizer4 5 6def main():7 path = os.path.expanduser("~/huggingface/Qwen3-0.6B/")8 tokenizer = AutoTokenizer.from_pretrained(path)9 llm = LLM(path, enforce_eager=True, tensor_parallel_size=1)10 11 sampling_params = SamplingParams(temperature=0.6, max_tokens=256)12 prompts = [13 "introduce yourself",14 "list all prime numbers within 100",15 ]16 prompts = [17 tokenizer.apply_chat_template(18 [{"role": "user", "content": prompt}],19 tokenize=False,20 add_generation_prompt=True,21 )22 for prompt in prompts23 ]24 outputs = llm.generate(prompts, sampling_params)25 26 for prompt, output in zip(prompts, outputs):27 print("\n")28 print(f"Prompt: {prompt!r}")29 print(f"Completion: {output['text']!r}")30 31 32if __name__ == "__main__":33 main()1.3bench.py: the throughput benchmark
- 256 requests with prompts of 100–1024 random token ids drawn from
[0, 10000], andmax_tokensdrawn from 100–1024.seed(0)makes the workload repeatable. ignore_eos=Truemakes every request produce exactlymax_tokenstokens, so the total output token count is known in advance.enforce_eager=Falseturns CUDA graphs on;max_model_len=4096covers the longest request (1024 + 1024).- A one-prompt
generate(["Benchmark: "], SamplingParams())runs before the timer starts, so first-call costs stay out of the measurement. - Throughput is output tokens divided by wall time, and the wall time includes all prefill work.
- The commented lines switch the same script to vLLM, which takes token-id prompts as
dict(prompt_token_ids=p).
1import os2import time3from random import randint, seed4from nanovllm import LLM, SamplingParams5# from vllm import LLM, SamplingParams6 7 8def main():9 seed(0)10 num_seqs = 25611 max_input_len = 102412 max_ouput_len = 102413 14 path = os.path.expanduser("~/huggingface/Qwen3-0.6B/")15 llm = LLM(path, enforce_eager=False, max_model_len=4096)16 17 prompt_token_ids = [[randint(0, 10000) for _ in range(randint(100, max_input_len))] for _ in range(num_seqs)]18 sampling_params = [SamplingParams(temperature=0.6, ignore_eos=True, max_tokens=randint(100, max_ouput_len)) for _ in range(num_seqs)]19 # uncomment the following line for vllm20 # prompt_token_ids = [dict(prompt_token_ids=p) for p in prompt_token_ids]21 22 llm.generate(["Benchmark: "], SamplingParams())23 t = time.time()24 llm.generate(prompt_token_ids, sampling_params, use_tqdm=False)25 t = (time.time() - t)26 total_tokens = sum(sp.max_tokens for sp in sampling_params)27 throughput = total_tokens / t28 print(f"Total: {total_tokens}tok, Time: {t:.2f}s, Throughput: {throughput:.2f}tok/s")29 30 31if __name__ == "__main__":32 main()The README reports this run on an RTX 4070 Laptop GPU (8 GB) with Qwen3-0.6B: vLLM produced 133,966 tokens in 98.37 s (1361.84 tok/s), nano-vllm the same tokens in 93.41 s (1434.13 tok/s).
1.4Packaging and dependencies
Python 3.10–3.12. Five direct dependencies, each with one job in the code:
| Dependency | Used for | Where |
|---|---|---|
torch>=2.4.0 | tensors, NCCL, CUDA graphs, torch.compile, spawn-based multiprocessing | everywhere |
triton>=3.0.0 | the KV-cache store kernel; also the backend torch.compile generates kernels with | layers/attention.py |
transformers>=4.51.0 | AutoConfig, AutoTokenizer, the Qwen3Config type | config.py, llm_engine.py, models/qwen3.py |
flash-attn | flash_attn_varlen_func (prefill) and flash_attn_with_kvcache (decode) | layers/attention.py |
xxhash | 64-bit hashes of KV blocks for prefix caching | engine/block_manager.py |
safetensors, numpy and tqdm are imported too; they arrive transitively through transformers.
1[build-system]2requires = ["setuptools>=61"]3build-backend = "setuptools.build_meta"4 5[project]6name = "nano-vllm"7version = "0.2.0"8authors = [{ name = "Xingkai Yu" }]9license = "MIT"10license-files = ["LICENSE"]11readme = "README.md"12description = "a lightweight vLLM implementation built from scratch"13requires-python = ">=3.10,<3.13"14dependencies = [15 "torch>=2.4.0",16 "triton>=3.0.0",17 "transformers>=4.51.0",18 "flash-attn",19 "xxhash",20]21 22[project.urls]23Homepage="https://github.com/GeeeekExplorer/nano-vllm"24 25[tool.setuptools.packages.find]26where = ["."]27include = ["nanovllm*"]Step 02Configuration
Two small dataclasses carry every setting. Config is per engine; SamplingParams is per request.
2.1Config: every engine knob
Config is a @dataclass(slots=True). slots=True fixes the attribute set and makes attribute access a little cheaper. Three fields start as placeholders and are filled later by other objects: hf_config in __post_init__, eos by LLMEngine, and num_kvcache_blocks by ModelRunner.allocate_kv_cache.
1import os2from dataclasses import dataclass3from transformers import AutoConfig4 5 6@dataclass(slots=True)7class Config:8 model: str9 max_num_batched_tokens: int = 1638410 max_num_seqs: int = 51211 max_model_len: int = 409612 gpu_memory_utilization: float = 0.913 tensor_parallel_size: int = 114 enforce_eager: bool = False15 hf_config: AutoConfig | None = None16 eos: int = -117 kvcache_block_size: int = 25618 num_kvcache_blocks: int = -119 20 def __post_init__(self):21 assert os.path.isdir(self.model)22 assert self.kvcache_block_size % 256 == 023 assert 1 <= self.tensor_parallel_size <= 824 self.hf_config = AutoConfig.from_pretrained(self.model)25 self.max_model_len = min(self.max_model_len, self.hf_config.max_position_embeddings)| Field | Default | Read by | Effect |
|---|---|---|---|
model | required | Config, LLMEngine, loader | Local checkpoint directory. Source of config.json, the tokenizer and *.safetensors. |
max_num_batched_tokens | 16384 | Scheduler, warmup | Token budget of one prefill step. A prompt longer than this is split into chunks. Also sizes the warmup batch. |
max_num_seqs | 512 | Scheduler, warmup, CUDA graphs | Most sequences in one step, prefill or decode. Largest captured graph is min(max_num_seqs, 512). |
max_model_len | 4096 | warmup, CUDA graphs | Clamped to the model's max_position_embeddings (40,960 for Qwen3-0.6B). Sets the warmup sequence length and the width of the graph's block-table buffer. Requests are not checked against it (Step 14). |
gpu_memory_utilization | 0.9 | allocate_kv_cache | Fraction of total GPU memory the process may plan for. Whatever weights, activations and other processes leave inside it becomes KV cache. |
tensor_parallel_size | 1 | LLMEngine, ModelRunner, all TP layers | Number of GPU processes. Asserted to be 1–8. |
enforce_eager | False | ModelRunner | True skips CUDA graph capture; every forward runs eagerly. |
hf_config | None | ModelRunner, model | Set to AutoConfig.from_pretrained(model) in __post_init__. |
eos | −1 | Scheduler | Set by LLMEngine to tokenizer.eos_token_id. |
kvcache_block_size | 256 | everything that touches blocks | Tokens per KV block. Must be a multiple of 256, the page size FlashAttention's paged-KV kernels require. |
num_kvcache_blocks | −1 | Scheduler → BlockManager | Set by allocate_kv_cache from measured free memory. |
2.2SamplingParams: per-request decoding
Three fields. The assertion forbids temperature 0: the sampler divides logits by the temperature (Step 10.7), so greedy decoding is not available. There is no top-k, top-p or repetition penalty. Sequence.__init__ copies the three values onto the sequence, so the params object is not referenced afterwards.
1from dataclasses import dataclass2 3 4@dataclass(slots=True)5class SamplingParams:6 temperature: float = 1.07 max_tokens: int = 648 ignore_eos: bool = False9 10 def __post_init__(self):11 assert self.temperature > 1e-10, "greedy sampling is not permitted"Step 03The engine loop: LLMEngine
LLMEngine owns the tokenizer, the scheduler and rank 0's model runner, and spawns the other ranks. Its generate method is a loop around step(), and each step is one schedule, one forward, one round of bookkeeping.
3.1Construction, and why the order matters
LLM(path, …) does before it can generate. Ranks meet twice: inside init_process_group and at the shared-memory barrier. The highlighted path is the one ordering constraint that shapes LLMEngine.__init__: the scheduler can only be built after rank 0 has measured how many KV blocks fit, because BlockManager needs that number.- Filter kwargs into
Config.dataclasses.fields(Config)gives the allowed names; everything else is dropped. - Set
Sequence.block_size. It is a class attribute, so everySequencein this process uses the configured block size fornum_blocksand friends. The assignment happens only in the main process (see Step 14 for what that means for workers). - Spawn ranks 1…N−1.
mp.get_context("spawn")starts fresh interpreters, which CUDA requires; a forked child cannot reuse the parent's CUDA context. The process target is theModelRunnerclass itself, so constructing it is the whole life of the worker: its__init__ends inself.loop(). Each worker gets oneEvent; rank 0 later receives the list of all of them. - Build rank 0's
ModelRunnerin this process. It must come after the spawns becauseinit_process_groupblocks until all ranks have joined. This call loads the model, measures memory, allocates the KV cache and captures CUDA graphs (Step 07). - Load the tokenizer and record EOS.
config.eosis set after the workers were spawned, so their copies ofconfigkeep −1. They never need it. - Build the
Schedulerlast. ItsBlockManagerneedsconfig.num_kvcache_blocks, which only exists after step 4. Spawned workers received a pickled copy ofconfigearlier, so this value is rank 0's own measurement. - Register
exitwithatexit, so workers are told to stop when the interpreter exits.
1import atexit2from dataclasses import fields3from time import perf_counter4from tqdm.auto import tqdm5from transformers import AutoTokenizer6import torch.multiprocessing as mp7 8from nanovllm.config import Config9from nanovllm.sampling_params import SamplingParams10from nanovllm.engine.sequence import Sequence11from nanovllm.engine.scheduler import Scheduler12from nanovllm.engine.model_runner import ModelRunner13 14 15class LLMEngine:16 17 def __init__(self, model, **kwargs):18 config_fields = {field.name for field in fields(Config)}19 config_kwargs = {k: v for k, v in kwargs.items() if k in config_fields}20 config = Config(model, **config_kwargs)21 Sequence.block_size = config.kvcache_block_size22 self.ps = []23 self.events = []24 ctx = mp.get_context("spawn")25 for i in range(1, config.tensor_parallel_size):26 event = ctx.Event()27 process = ctx.Process(target=ModelRunner, args=(config, i, event))28 process.start()29 self.ps.append(process)30 self.events.append(event)31 self.model_runner = ModelRunner(config, 0, self.events)32 self.tokenizer = AutoTokenizer.from_pretrained(config.model, use_fast=True)33 config.eos = self.tokenizer.eos_token_id34 self.scheduler = Scheduler(config)35 atexit.register(self.exit)3.2add_request: text or token ids in, a Sequence out
Strings go through tokenizer.encode; lists of ints are used as they are, which is how bench.py sends random prompts. The new Sequence is appended to scheduler.waiting; nothing else happens until the next step().
43 def add_request(self, prompt: str | list[int], sampling_params: SamplingParams):44 if isinstance(prompt, str):45 prompt = self.tokenizer.encode(prompt)46 seq = Sequence(prompt, sampling_params)47 self.scheduler.add(seq)3.3step: one iteration of the serving loop
scheduler.schedule()returns the batch and whether it is a prefill batch.num_tokensis measured before the forward, becausepostprocessresets everynum_scheduled_tokensto 0. The sign carries the phase: positive means “prefill tokens processed”, negative means “number of sequences decoded”.generateuses the sign to pick which throughput figure to update.model_runner.call("run", seqs, is_prefill)runs the forward on every rank and returns one token id per sequence on rank 0.scheduler.postprocessappends tokens, registers newly full blocks for prefix caching, and retires finished sequences.- The return value lists
(seq_id, completion_token_ids)only for sequences that finished in this step.
49 def step(self):50 seqs, is_prefill = self.scheduler.schedule()51 num_tokens = sum(seq.num_scheduled_tokens for seq in seqs) if is_prefill else -len(seqs)52 token_ids = self.model_runner.call("run", seqs, is_prefill)53 self.scheduler.postprocess(seqs, token_ids, is_prefill)54 outputs = [(seq.seq_id, seq.completion_token_ids) for seq in seqs if seq.is_finished]55 return outputs, num_tokens56 57 def is_finished(self):58 return self.scheduler.is_finished()3.4generate: batch in, ordered results out
- A single
SamplingParamsis repeated for every prompt; a list is used element-wise. - All requests are enqueued before the first step, so the scheduler sees the whole workload at once.
- The progress bar counts finished requests. Its postfix shows the most recent prefill rate (tokens / step time) and decode rate (sequences / step time, one token each).
- Results are collected in a dict keyed by
seq_idand then sorted.seq_idcomes from a process-wideitertools.count(), so sorting restores submission order even though requests finish in any order, and it keeps working across repeatedgeneratecalls. tokenizer.decode(token_ids)does not skip special tokens. A request that stopped on EOS includes the EOS token in bothtoken_idsandtext(for Qwen3 chat models,<|im_end|>).- The return annotation says
list[str]; the function actually returns a list of dicts.
57 def is_finished(self):58 return self.scheduler.is_finished()59 60 def generate(61 self,62 prompts: list[str] | list[list[int]],63 sampling_params: SamplingParams | list[SamplingParams],64 use_tqdm: bool = True,65 ) -> list[str]:66 pbar = tqdm(total=len(prompts), desc="Generating", dynamic_ncols=True, disable=not use_tqdm)67 if not isinstance(sampling_params, list):68 sampling_params = [sampling_params] * len(prompts)69 for prompt, sp in zip(prompts, sampling_params):70 self.add_request(prompt, sp)71 outputs = {}72 prefill_throughput = decode_throughput = 0.73 while not self.is_finished():74 t = perf_counter()75 output, num_tokens = self.step()76 if num_tokens > 0:77 prefill_throughput = num_tokens / (perf_counter() - t)78 else:79 decode_throughput = -num_tokens / (perf_counter() - t)80 pbar.set_postfix({81 "Prefill": f"{int(prefill_throughput)}tok/s",82 "Decode": f"{int(decode_throughput)}tok/s",83 })84 for seq_id, token_ids in output:85 outputs[seq_id] = token_ids86 pbar.update(1)87 pbar.close()88 outputs = [outputs[seq_id] for seq_id in sorted(outputs.keys())]89 outputs = [{"text": self.tokenizer.decode(token_ids), "token_ids": token_ids} for token_ids in outputs]90 return outputs3.5exit: shutting down the ranks
call("exit") writes "exit" into shared memory so every worker leaves its loop and runs its own ModelRunner.exit, then runs rank 0's exit locally. Deleting the runner and joining the processes finishes the shutdown.
37 def exit(self):38 self.model_runner.call("exit")39 del self.model_runner40 for p in self.ps:41 p.join()3.6The life of one request
generate(). Everything inside the frame is one step(): exactly one scheduling decision, one forward pass on every rank, one round of bookkeeping. A step is either all-prefill or all-decode, never a mix.Keep this picture in mind for the next four steps. Steps 04–06 explain the three left-hand lifelines. Steps 07–11 explain the two on the right.
Step 04The unit of work: Sequence
A Sequence is one request: its token ids, four counters that say how far its KV cache has progressed, its block table, and its sampling settings. Almost every decision in the scheduler and runner is a comparison between these counters.
4.1Identity, counters and block table
1from copy import copy2from enum import Enum, auto3from itertools import count4 5from nanovllm.sampling_params import SamplingParams6 7 8class SequenceStatus(Enum):9 WAITING = auto()10 RUNNING = auto()11 FINISHED = auto()12 13 14class Sequence:15 block_size = 25616 counter = count()17 18 def __init__(self, token_ids: list[int], sampling_params = SamplingParams()):19 self.seq_id = next(Sequence.counter)20 self.status = SequenceStatus.WAITING21 self.token_ids = copy(token_ids)22 self.last_token = token_ids[-1]23 self.num_tokens = len(self.token_ids)24 self.num_prompt_tokens = len(token_ids)25 self.num_cached_tokens = 026 self.num_scheduled_tokens = 027 self.is_prefill = True28 self.block_table = []29 self.temperature = sampling_params.temperature30 self.max_tokens = sampling_params.max_tokens31 self.ignore_eos = sampling_params.ignore_eos- seq_id
- From the class-level
itertools.count(), shared by every sequence in the process. Used only to order outputs. - token_ids
- Prompt plus generated tokens. The constructor copies the caller's list so
append_tokennever mutates it. - num_tokens
len(token_ids), kept as a counter because worker processes receive aSequencewithout its token list during decode (4.4).- num_prompt_tokens
- Fixed at creation. Splits
token_idsinto prompt and completion. - num_cached_tokens
- How many leading tokens already have their K/V in the cache: prefix-cache hits, earlier prefill chunks and every decoded token. Advanced only in
postprocess; reset to 0 bydeallocate. - num_scheduled_tokens
- How many tokens this step will compute, starting at position
num_cached_tokens. Set by the scheduler, read byprepare_*andhash_blocks, reset to 0 inpostprocess. - is_prefill
- True until the sequence is first scheduled for decode, and again after a preemption. It only decides how the sequence is pickled.
- block_table
- Physical KV block ids, one per logical block, filled by
BlockManager.
Sequence, drawn with a toy block size of 4 (real blocks hold 256 tokens). Cells are token positions. In (b) the newest token, position 16, has been appended by postprocess but its K/V is not in the cache yet, so in decode num_cached_tokens is always num_tokens − 1. Position 16 is the first token of a fifth block, which is why may_append added block 5.4.2Block arithmetic and list behaviour
__len__and__getitem__let a sequence behave like its token list, so the runner writeslen(seq)andseq[start:end].num_blocksisceil(num_tokens / block_size), written with integer arithmetic.last_block_num_tokensis how many tokens occupy the last block, from 1 toblock_size. Decode uses it to find the slot of the newest token.block(i)returns the token ids of logical blocki; the last block may be partial. Prefix caching hashes these lists.append_tokenupdatestoken_ids,last_tokenandnum_tokenstogether.
33 def __len__(self):34 return self.num_tokens35 36 def __getitem__(self, key):37 return self.token_ids[key]38 39 @property40 def is_finished(self):41 return self.status == SequenceStatus.FINISHED42 43 @property44 def num_completion_tokens(self):45 return self.num_tokens - self.num_prompt_tokens46 47 @property48 def prompt_token_ids(self):49 return self.token_ids[:self.num_prompt_tokens]50 51 @property52 def completion_token_ids(self):53 return self.token_ids[self.num_prompt_tokens:]54 55 @property56 def num_blocks(self):57 return (self.num_tokens + self.block_size - 1) // self.block_size58 59 @property60 def last_block_num_tokens(self):61 return self.num_tokens - (self.num_blocks - 1) * self.block_size62 63 def block(self, i):64 assert 0 <= i < self.num_blocks65 return self.token_ids[i*self.block_size: (i+1)*self.block_size]66 67 def append_token(self, token_id: int):68 self.token_ids.append(token_id)69 self.last_token = token_id70 self.num_tokens += 14.3Status transitions
8class SequenceStatus(Enum):9 WAITING = auto()10 RUNNING = auto()11 FINISHED = auto()SequenceStatus has three values. A sequence becomes RUNNING inside schedule() as soon as its final prefill chunk is put in the batch, so a request whose first sampled token is EOS goes WAITING → RUNNING → FINISHED within a single step.4.4Compact pickling for worker ranks
With tensor parallelism, rank 0 sends the batch's sequences to the other ranks through shared memory every step (Step 07.2). __getstate__ keeps that message small by sending a 6-tuple instead of the object's __dict__:
- the four counters and the block table, which the slot and position computations need;
- the whole token list while the sequence is in prefill, because
prepare_prefillslicesseq[start:end]; - only
last_tokenin decode, becauseprepare_decodeneeds nothing else. A 4,000-token sequence costs one int instead of 4,000.
Nothing else crosses: no seq_id, status, temperature, max_tokens or ignore_eos. Workers never sample and never schedule, so they never need them.
72 def __getstate__(self):73 last_state = self.last_token if not self.is_prefill else self.token_ids74 return (self.num_tokens, self.num_prompt_tokens, self.num_cached_tokens, self.num_scheduled_tokens, self.block_table, last_state)75 76 def __setstate__(self, state):77 self.num_tokens, self.num_prompt_tokens, self.num_cached_tokens, self.num_scheduled_tokens, self.block_table, last_state = state78 if isinstance(last_state, list):79 self.token_ids = last_state80 self.last_token = self.token_ids[-1]81 else:82 self.token_ids = []83 self.last_token = last_stateStep 05Paged KV memory: BlockManager
BlockManager is a CPU-side allocator for GPU memory it never touches. It hands out block ids, counts how many sequences share each block, and remembers which token prefixes each block holds so later requests can reuse them.
5.1Blocks, block tables and slots
block_table is the only link between the two. A slot is a row index into the flattened pool, slot = block_id × block_size + offset; the Triton store kernel writes a token's K and V at that slot, and FlashAttention reads them back through the block table.The physical pool is one big tensor allocated by the runner (Step 07.4). The block manager only deals in integer ids from 0 to num_kvcache_blocks − 1. Two numbers follow from a block table and a position: the block block_table[pos // block_size] and the slot block_id × block_size + pos % block_size. The runner computes slots; the block manager decides which ids appear in the table.
5.2Block and the four bookkeeping structures
A Block records how many sequences use it (ref_count) and, once it is full and computed, the hash of the prefix it completes and its token ids. reset() sets ref_count to 1, not 0, because it is only called on a block that is being handed to a sequence.
1from collections import deque2import xxhash3import numpy as np4 5from nanovllm.engine.sequence import Sequence6 7 8class Block:9 10 def __init__(self, block_id):11 self.block_id = block_id12 self.ref_count = 013 self.hash = -114 self.token_ids = []15 16 def update(self, hash: int, token_ids: list[int]):17 self.hash = hash18 self.token_ids = token_ids19 20 def reset(self):21 self.ref_count = 122 self.hash = -123 self.token_ids = []24 25 26class BlockManager:27 28 def __init__(self, num_blocks: int, block_size: int):29 self.block_size = block_size30 self.blocks: list[Block] = [Block(i) for i in range(num_blocks)]31 self.hash_to_block_id: dict[int, int] = dict()32 self.free_block_ids: deque[int] = deque(range(num_blocks))33 self.used_block_ids: set[int] = set()| Structure | Type | Meaning |
|---|---|---|
blocks | list[Block] | All block objects, indexed by id. |
hash_to_block_id | dict[int, int] | Prefix hash → the block holding that prefix's last block of K/V. This is the prefix cache. |
free_block_ids | deque[int] | Blocks with ref_count == 0. Allocation pops from the left; freeing appends to the right, so the least recently freed block is recycled first. |
used_block_ids | set[int] | Blocks with ref_count > 0. Every id is in exactly one of the free deque and this set. |
5.3Chained block hashes
compute_hash feeds xxh64 the previous block's hash (8 little-endian bytes) followed by the block's token ids as a NumPy array. Chaining makes a block's hash identify the entire prefix up to and including that block. That is the right key: in every layer after the first, a token's K and V depend on every earlier token, so the same four tokens after a different prefix have different K/V.
35 @classmethod36 def compute_hash(cls, token_ids: list[int], prefix: int = -1):37 h = xxhash.xxh64()38 if prefix != -1:39 h.update(prefix.to_bytes(8, "little"))40 h.update(np.array(token_ids).tobytes())41 return h.intdigest()hash_to_block_id, confirms the stored token ids are equal, and shares blocks 3 and 7 (their ref_count goes to 2). Its last block is never looked up, so at least that block is always computed and produces the logits for B's first output token.5.4can_allocate: count cache hits, check capacity
The scheduler calls this before admitting a waiting sequence that has no blocks yet. It changes nothing; it answers how many leading blocks are cache hits, or −1 if the sequence cannot fit.
- Walk the logical blocks, recomputing the hash chain. Stop at the first block whose hash is unknown or whose stored token ids differ. The token comparison protects against a 64-bit hash collision.
- The loop runs to
num_blocks − 1, so the last block is never matched, even when the prompt is block-aligned and fully cached. At least one block's worth of tokens is always computed, and its final position produces the logits for the next token. num_new_blocksstarts at the full block count. A hit on a block that is in use by another sequence costs no new memory, so it is subtracted. A hit on a block sitting in the free deque still has to be taken out of the deque, so it still counts.- Return −1 if the free deque is shorter than
num_new_blocks, otherwise the number of cached blocks.
58 def can_allocate(self, seq: Sequence) -> int:59 h = -160 num_cached_blocks = 061 num_new_blocks = seq.num_blocks62 for i in range(seq.num_blocks - 1):63 token_ids = seq.block(i)64 h = self.compute_hash(token_ids, h)65 block_id = self.hash_to_block_id.get(h, -1)66 if block_id == -1 or self.blocks[block_id].token_ids != token_ids:67 break68 num_cached_blocks += 169 if block_id in self.used_block_ids:70 num_new_blocks -= 171 if len(self.free_block_ids) < num_new_blocks:72 return -173 return num_cached_blocks5.5allocate: build the block table
- For each cached block, recompute the hash and look it up. If the block is in use, increment its
ref_count. If it is free, bring it back:ref_count = 1, remove it from the free deque (an O(n)deque.remove), add it to the used set. Its contents are untouched, so no recomputation is needed. - Every remaining logical block gets a fresh block from
_allocate_block. The whole sequence is allocated at once, even if chunked prefill will fill it over several steps. num_cached_tokensbecomesnum_cached_blocks × block_size; prefill will start computing at that position.
75 def allocate(self, seq: Sequence, num_cached_blocks: int):76 assert not seq.block_table77 h = -178 for i in range(num_cached_blocks):79 token_ids = seq.block(i)80 h = self.compute_hash(token_ids, h)81 block_id = self.hash_to_block_id[h]82 block = self.blocks[block_id]83 if block_id in self.used_block_ids:84 block.ref_count += 185 else:86 block.ref_count = 187 self.free_block_ids.remove(block_id)88 self.used_block_ids.add(block_id)89 seq.block_table.append(block_id)90 for i in range(num_cached_blocks, seq.num_blocks):91 seq.block_table.append(self._allocate_block())92 seq.num_cached_tokens = num_cached_blocks * self.block_size5.6_allocate_block: take a block, evict its old identity
Pop the leftmost free id. If the block still carries a hash and the cache map still points at this block, delete that map entry: the block is about to be overwritten. The == block_id check matters because one hash can be registered twice. Two identical prompts admitted in the same prefill step both get fresh blocks (neither is hashed yet), both are hashed in postprocess, and the map ends up pointing at the second. Recycling the first block later must not delete the entry that belongs to the second.
43 def _allocate_block(self) -> int:44 block_id = self.free_block_ids.popleft()45 block = self.blocks[block_id]46 assert block.ref_count == 047 if block.hash != -1 and self.hash_to_block_id.get(block.hash) == block_id:48 del self.hash_to_block_id[block.hash]49 block.reset()50 self.used_block_ids.add(block_id)51 return block_id5.7deallocate: release a sequence's blocks
Used when a sequence finishes and when it is preempted. Each block's ref_count drops by one; a block that reaches zero moves to the right end of the free deque, keeping its hash. The table is walked in reverse so a sequence's tail blocks enter the deque before its prefix blocks and are recycled first. Prefix blocks, the ones most likely to be shared again, survive longest.
53 def _deallocate_block(self, block_id: int):54 assert self.blocks[block_id].ref_count == 055 self.used_block_ids.remove(block_id)56 self.free_block_ids.append(block_id)94 def deallocate(self, seq: Sequence):95 for block_id in reversed(seq.block_table):96 block = self.blocks[block_id]97 block.ref_count -= 198 if block.ref_count == 0:99 self._deallocate_block(block_id)100 seq.num_cached_tokens = 0101 seq.block_table.clear()5.8can_append and may_append: growing during decode
When a decode step is scheduled, the token sampled in the previous step is already in token_ids but its K/V is not stored yet. This step stores it at position len(seq) − 1. That position starts a new logical block exactly when len(seq) % block_size == 1, and only then is a new block needed. can_append turns the comparison into 0 or 1 by letting Python treat the boolean as an int.
| len(seq) at schedule time | len % 4 | Position stored this step | Action (toy block size 4) |
|---|---|---|---|
| 8 | 0 | 7, last slot of block 1 | none; block 1 becomes full and is hashed in postprocess |
| 9 | 1 | 8, first slot of block 2 | may_append allocates block 2 |
| 10 | 2 | 9 | none |
| 11 | 3 | 10 | none |
| 12 | 0 | 11, block 2 full | none; block 2 hashed |
103 def can_append(self, seq: Sequence) -> bool:104 return len(self.free_block_ids) >= (len(seq) % self.block_size == 1)105 106 def may_append(self, seq: Sequence):107 if len(seq) % self.block_size == 1:108 seq.block_table.append(self._allocate_block())5.9hash_blocks: publish newly full blocks
Called from postprocess for every sequence in the batch, after the forward and before num_cached_tokens advances. start is the first block that was not yet full before this step; end is the number of full blocks after it. Every block in between is hashed, using the previous block's stored hash as the chain prefix, and registered in hash_to_block_id.
- It runs after decode steps as well as prefill. Blocks filled by generated tokens become cacheable, which is what makes recomputation after preemption cheap (Step 06.4).
- A partial last block is never hashed, so a block's hash always covers exactly
block_sizetokens. - Blocks are registered only after their K/V has really been written, so a later lookup can never hit a block whose contents are still pending.
110 def hash_blocks(self, seq: Sequence):111 start = seq.num_cached_tokens // self.block_size112 end = (seq.num_cached_tokens + seq.num_scheduled_tokens) // self.block_size113 if start == end: return114 h = self.blocks[seq.block_table[start - 1]].hash if start > 0 else -1115 for i in range(start, end):116 block = self.blocks[seq.block_table[i]]117 token_ids = seq.block(i)118 h = self.compute_hash(token_ids, h)119 block.update(h, token_ids)120 self.hash_to_block_id[h] = block.block_idStep 06Choosing each batch: Scheduler
The scheduler is 92 lines and holds two queues. schedule() picks the next batch, preempt() frees memory under pressure, and postprocess() applies the sampled tokens. Read it with Step 05 open: almost every branch is a call into the block manager.
6.1Queues and limits
Two deques. waiting holds new requests, preempted ones and a sequence whose prefill is part-way through its chunks. running holds sequences whose prompt is fully in the cache and that are generating. The limits and EOS id are copied from config, and the block manager is created here with rank 0's measured block count.
1from collections import deque2 3from nanovllm.config import Config4from nanovllm.engine.sequence import Sequence, SequenceStatus5from nanovllm.engine.block_manager import BlockManager6 7 8class Scheduler:9 10 def __init__(self, config: Config):11 self.max_num_seqs = config.max_num_seqs12 self.max_num_batched_tokens = config.max_num_batched_tokens13 self.eos = config.eos14 self.block_size = config.kvcache_block_size15 self.block_manager = BlockManager(config.num_kvcache_blocks, config.kvcache_block_size)16 self.waiting: deque[Sequence] = deque()17 self.running: deque[Sequence] = deque()18 19 def is_finished(self):20 return not self.waiting and not self.running21 22 def add(self, seq: Sequence):23 self.waiting.append(seq)6.2Prefill phase and chunked prefill
Prefill always goes first. If anything at all is scheduled in this loop, the step is a prefill step and no sequence decodes in it.
- First come, first served. The loop only looks at
waiting[0]. When the head cannot be served itbreaks; it never skips ahead to a smaller request. - Budget.
remainingis what is left ofmax_num_batched_tokens. A batch that has used it all stops. - Fresh or continuing. A sequence without blocks is new or was preempted:
can_allocatereports cache hits or −1, and the work is everything after the cached blocks. A sequence that already has blocks is mid-way through a chunked prefill, and its work is everything afternum_cached_tokens. - The chunking rule. If the work does not fit in
remainingand the batch already has a sequence, stop. Only the first sequence of a batch can be cut. That sequence always sees the full budget, so in practice a prompt is chunked only when its uncached part is longer thanmax_num_batched_tokens. With the defaults (16,384 tokens of budget,max_model_len4,096) that never happens; lower the budget or raisemax_model_lento see it. - Commit. Blocks are allocated only now that the sequence is definitely in the batch.
num_scheduled_tokensis the work, capped by the budget. - Done or not. If this chunk reaches
num_tokens, the sequence becomesRUNNINGand moves torunning. Otherwise it stays atwaiting[0]; it used the whole remaining budget, so the next iteration hitsremaining == 0and the loop ends.
25 def schedule(self) -> tuple[list[Sequence], bool]:26 scheduled_seqs = []27 num_batched_tokens = 028 29 # prefill30 while self.waiting and len(scheduled_seqs) < self.max_num_seqs:31 seq = self.waiting[0]32 remaining = self.max_num_batched_tokens - num_batched_tokens33 if remaining == 0:34 break35 if not seq.block_table:36 num_cached_blocks = self.block_manager.can_allocate(seq)37 if num_cached_blocks == -1:38 break39 num_tokens = seq.num_tokens - num_cached_blocks * self.block_size40 else:41 num_tokens = seq.num_tokens - seq.num_cached_tokens42 if remaining < num_tokens and scheduled_seqs: # only allow chunked prefill for the first seq43 break44 if not seq.block_table:45 self.block_manager.allocate(seq, num_cached_blocks)46 seq.num_scheduled_tokens = min(num_tokens, remaining)47 num_batched_tokens += seq.num_scheduled_tokens48 if seq.num_cached_tokens + seq.num_scheduled_tokens == seq.num_tokens:49 seq.status = SequenceStatus.RUNNING50 self.waiting.popleft()51 self.running.append(seq)52 scheduled_seqs.append(seq)53 54 if scheduled_seqs:55 return scheduled_seqs, True6.3Decode phase and preemption
- Only reached when nothing could be prefilled.
- Take sequences from the front of
running, up tomax_num_seqs. - Make sure the sequence can store one more token (
can_append). While it cannot, preempt the sequence at the back ofrunning, the most recently admitted one, which frees its blocks. If no other sequence is left, preempt this one and move on. - Python's
while … else: theelseblock runs when the loop condition turns false, meaning there is room, and is skipped after thebreakthat follows self-preemption. In theelsebranch the sequence is set to compute one token, markedis_prefill = Falseso it pickles compactly, and given a new block if it needs one. assert scheduled_seqsfires only if a single sequence needs more blocks than the whole pool, because then even preempting everything else frees nothing usable.running.extendleft(reversed(scheduled_seqs))puts the batch back at the front ofrunningin its original order.
57 # decode58 while self.running and len(scheduled_seqs) < self.max_num_seqs:59 seq = self.running.popleft()60 while not self.block_manager.can_append(seq):61 if self.running:62 self.preempt(self.running.pop())63 else:64 self.preempt(seq)65 break66 else:67 seq.num_scheduled_tokens = 168 seq.is_prefill = False69 self.block_manager.may_append(seq)70 scheduled_seqs.append(seq)71 assert scheduled_seqs72 self.running.extendleft(reversed(scheduled_seqs))73 return scheduled_seqs, False6.4preempt: give memory back, recompute later
The victim goes back to WAITING, is marked as prefill again, loses all its blocks and is put at the front of waiting, so it is the first sequence re-admitted. There is no swapping to CPU memory; the K/V will be recomputed. Two details make that cheaper than it sounds:
deallocatekeeps hashes, and decode steps hashed every block they filled. When the sequence is re-admitted,can_allocatefinds all of its full blocks that have not been recycled yet, and only the rest is recomputed.- The last token of a preempted sequence was sampled but never fed back. The re-prefill computes it, so the token sampled at the end of that prefill is a correct next token and
postprocessappends it.
75 def preempt(self, seq: Sequence):76 seq.status = SequenceStatus.WAITING77 seq.is_prefill = True78 self.block_manager.deallocate(seq)79 self.waiting.appendleft(seq)6.5postprocess: apply the sampled tokens
hash_blocksfirst, whilenum_cached_tokensandnum_scheduled_tokensstill describe what was just computed.- Advance
num_cached_tokens, clearnum_scheduled_tokens. - A prefill chunk that did not reach the end of the sequence produced a token for a position in the middle of the prompt. It is discarded with
continue. - Otherwise append the token. Finish when the token is EOS (unless
ignore_eos) or the completion has reachedmax_tokens. The token is appended before the check, so a finishing EOS is part of the output. - A finished sequence frees its blocks immediately and leaves
running(an O(n)deque.remove).
81 def postprocess(self, seqs: list[Sequence], token_ids: list[int], is_prefill: bool):82 for seq, token_id in zip(seqs, token_ids):83 self.block_manager.hash_blocks(seq)84 seq.num_cached_tokens += seq.num_scheduled_tokens85 seq.num_scheduled_tokens = 086 if is_prefill and seq.num_cached_tokens < seq.num_tokens:87 continue88 seq.append_token(token_id)89 if (not seq.ignore_eos and token_id == self.eos) or seq.num_completion_tokens == seq.max_tokens:90 seq.status = SequenceStatus.FINISHED91 self.block_manager.deallocate(seq)92 self.running.remove(seq)6.6Simulator: run the scheduler yourself
This runs a line-by-line JavaScript port of sequence.py, block_manager.py and scheduler.py, plus the slot arithmetic of prepare_prefill and prepare_decode. Block size is 4 instead of 256 so a whole sequence fits on screen, a 32-bit FNV hash stands in for xxh64, and the model is replaced by random tokens 1–9 (no EOS, so sequences finish at max_tokens). Everything else follows the Python.
With the default settings, stepping through shows each mechanism in turn: step 1 admits A alone, step 2 admits B with a two-block prefix hit on A's blocks, steps 3–4 chunk C's 22-token prompt across the 16-token budget, step 7 preempts C when the pool runs dry, and step 9 re-admits C with six cached blocks and recomputes one token.
Things to try: set num_kvcache_blocks to 8 and max_tokens to 12, then run until idle; C alone outgrows the pool and the scheduler's assertion fires, the same failure Step 14 describes. After everything finishes, add A again: its blocks are still in the free list with their hashes, so it is admitted with a prefix hit.
Step 07From sequences to tensors: ModelRunner
One ModelRunner exists per GPU. It builds the model, measures memory, owns the KV cache, captures CUDA graphs, converts a list of Sequences into the flat tensors FlashAttention wants, and, on rank 0, samples. On other ranks it is a loop that waits for instructions.
7.1Boot sequence
- Join the process group. NCCL, rendezvous at
tcp://localhost:2333, so all ranks must be on one machine.torch.cuda.set_device(rank)pins rank i to the i-th visible GPU. - Build the model directly on the GPU. The default dtype becomes the checkpoint's (
hf_config.dtype, bfloat16 for Qwen3) and the default device becomes CUDA, so everytorch.emptyinside the layer constructors allocates GPU memory in model precision. Nothing is initialised on the CPU first. Older transformers releases call this attributetorch_dtype; the code switched todtypein commit8d63a98, and an install whose config objects lackdtypefails on this line. - Load weights with
load_model(Step 11) and create theSampler. - Warm up before allocating the KV cache, so the measured activation peak does not include it (7.3).
- Allocate the KV cache (7.4), then capture CUDA graphs unless
enforce_eager(7.9). Capture has to come after the cache exists, because the captured kernels hold the cache's addresses. - Restore the defaults to CPU and the previous dtype. From here on, tensors built from Python lists in
prepare_*start on the CPU and are copied over explicitly. - Set up shared memory when
world_size > 1. Rank 0 creates a 1 MiB segment named"nanovllm"and then waits at the barrier; workers wait at the barrier first, so the segment exists before they attach, then enterloop()and stay there.
1import pickle2import torch3import torch.distributed as dist4from multiprocessing.synchronize import Event5from multiprocessing.shared_memory import SharedMemory6 7from nanovllm.config import Config8from nanovllm.engine.sequence import Sequence9from nanovllm.models.qwen3 import Qwen3ForCausalLM10from nanovllm.layers.sampler import Sampler11from nanovllm.utils.context import set_context, get_context, reset_context12from nanovllm.utils.loader import load_model13 14 15class ModelRunner:16 17 def __init__(self, config: Config, rank: int, event: Event | list[Event]):18 self.config = config19 hf_config = config.hf_config20 self.block_size = config.kvcache_block_size21 self.enforce_eager = config.enforce_eager22 self.world_size = config.tensor_parallel_size23 self.rank = rank24 self.event = event25 26 dist.init_process_group("nccl", "tcp://localhost:2333", world_size=self.world_size, rank=rank)27 torch.cuda.set_device(rank)28 default_dtype = torch.get_default_dtype()29 torch.set_default_dtype(hf_config.dtype)30 torch.set_default_device("cuda")31 self.model = Qwen3ForCausalLM(hf_config)32 load_model(self.model, config.model)33 self.sampler = Sampler()34 self.warmup_model()35 self.allocate_kv_cache()36 if not self.enforce_eager:37 self.capture_cudagraph()38 torch.set_default_device("cpu")39 torch.set_default_dtype(default_dtype)40 41 if self.world_size > 1:42 if rank == 0:43 self.shm = SharedMemory(name="nanovllm", create=True, size=2**20)44 dist.barrier()45 else:46 dist.barrier()47 self.shm = SharedMemory(name="nanovllm")48 self.loop()7.2The control plane: call, write_shm, read_shm, loop
Sequence objects, encoded compactly by __getstate__. Activations never touch shared memory; they are combined by NCCL collectives inside the forward pass, which also keeps rank 0 from racing ahead of the workers.call(name, *args)is the only entry point the engine uses. On rank 0 with several ranks it first broadcasts the call, then runs the method locally by name throughgetattr. In practice only"run"and"exit"go through it.- The message format is a 4-byte little-endian length followed by
pickle.dumps([name, *args]). The segment is fixed at 220 bytes, and the payload has to fit. - Each worker has its own
Event. A worker waits on it, reads, clears it, then runs the method. - No acknowledgement flows back. Ordering still holds: every forward pass contains NCCL collectives, so rank 0 cannot finish a
rununtil every worker has entered the samerun, which they only do after reading and clearing. Rank 0 therefore never overwrites a message a worker has not read yet.
61 def loop(self):62 while True:63 method_name, args = self.read_shm()64 self.call(method_name, *args)65 if method_name == "exit":66 break67 68 def read_shm(self):69 assert self.world_size > 1 and self.rank > 070 self.event.wait()71 n = int.from_bytes(self.shm.buf[0:4], "little")72 method_name, *args = pickle.loads(self.shm.buf[4:n+4])73 self.event.clear()74 return method_name, args75 76 def write_shm(self, method_name, *args):77 assert self.world_size > 1 and self.rank == 078 data = pickle.dumps([method_name, *args])79 n = len(data)80 self.shm.buf[0:4] = n.to_bytes(4, "little")81 self.shm.buf[4:n+4] = data82 for event in self.event:83 event.set()84 85 def call(self, method_name, *args):86 if self.world_size > 1 and self.rank == 0:87 self.write_shm(method_name, *args)88 method = getattr(self, method_name, None)89 return method(*args)7.3warmup_model: measure the worst prefill
The warmup runs the largest prefill the scheduler can produce, so the activation peak it records is an upper bound for real steps.
- Reset the allocator cache and peak statistics.
seq_len = min(max_num_batched_tokens, max_model_len)andnum_seqs = min(max_num_batched_tokens // seq_len, max_num_seqs). With the defaults that is 4 dummy sequences of 4,096 zero tokens, 16,384 tokens in total.- The dummy sequences have no block table.
prepare_prefillskips their slot mapping, andAttentionskips the KV store becausek_cacheis still the empty placeholder. - The first call of each
@torch.compilefunction happens here, so compilation cost is paid at start-up. - Every rank runs its own warmup inside its own constructor. The collectives inside the forward keep them in lock-step.
91 def warmup_model(self):92 torch.cuda.empty_cache()93 torch.cuda.reset_peak_memory_stats()94 max_num_batched_tokens, max_model_len = self.config.max_num_batched_tokens, self.config.max_model_len95 seq_len = min(max_num_batched_tokens, max_model_len)96 num_seqs = min(max_num_batched_tokens // seq_len, self.config.max_num_seqs)97 seqs = [Sequence([0] * seq_len) for _ in range(num_seqs)]98 for seq in seqs:99 seq.num_scheduled_tokens = seq_len100 self.run(seqs, True)101 torch.cuda.empty_cache()7.4allocate_kv_cache: size the pool from what is left
Four memory numbers are read right after warmup:
- total, free
torch.cuda.mem_get_info(), device-wide.used = total − freeincludes the CUDA context, other processes on the GPU and everything PyTorch holds.- current
- Bytes PyTorch has allocated right now: essentially the weights.
- peak
- The highest allocation during warmup: weights plus the worst activations.
The KV budget is total × gpu_memory_utilization − used − (peak − current): the allowed share, minus what is in use, minus room to bring the activation peak back. Dividing by the size of one block gives num_kvcache_blocks, written into config for the scheduler to read.
One block stores K and V (the factor 2) for every layer, block_size tokens, and this rank's KV heads:
block_bytes = 2 × num_hidden_layers × block_size × (num_key_value_heads / tp) × head_dim × dtype.itemsize
For Qwen3-0.6B at tp = 1: 2 × 28 × 256 × 8 × 128 × 2 = 29,360,128 bytes, exactly 28 MiB per block, or 112 KiB per token.
103 def allocate_kv_cache(self):104 config = self.config105 hf_config = config.hf_config106 free, total = torch.cuda.mem_get_info()107 used = total - free108 peak = torch.cuda.memory_stats()["allocated_bytes.all.peak"]109 current = torch.cuda.memory_stats()["allocated_bytes.all.current"]110 num_kv_heads = hf_config.num_key_value_heads // self.world_size111 head_dim = getattr(hf_config, "head_dim", hf_config.hidden_size // hf_config.num_attention_heads)112 block_bytes = 2 * hf_config.num_hidden_layers * self.block_size * num_kv_heads * head_dim * hf_config.dtype.itemsize113 config.num_kvcache_blocks = int(total * config.gpu_memory_utilization - used - peak + current) // block_bytes114 assert config.num_kvcache_blocks > 0115 self.kv_cache = torch.empty(2, hf_config.num_hidden_layers, config.num_kvcache_blocks, self.block_size, num_kv_heads, head_dim)116 layer_id = 0117 for module in self.model.modules():118 if hasattr(module, "k_cache") and hasattr(module, "v_cache"):119 module.k_cache = self.kv_cache[0, layer_id]120 module.v_cache = self.kv_cache[1, layer_id]121 layer_id += 1allocate_kv_cache walks model.modules() and hands every Attention layer a view of its own slice, so writes through k_cache land in the shared tensor. Columns in the zoomed panel are offsets inside a block (8 drawn, 256 real); a slot is one row of kv_heads × head_dim contiguous values, which is why the store kernel can write a token with one vector load and one vector store.Weights in the calculator are computed from each model's config.json (bf16, divided by tp except the norms). Qwen3-0.6B comes to 596,049,920 parameters, about 1.11 GiB.
7.5prepare_prefill: pack a varlen batch
cu_seqlens_q marks where each sequence's new tokens start; cu_seqlens_k does the same for the keys each sequence attends to, which include its cached prefix (green). slot_mapping tells the store kernel where each new token's K/V goes. Order matters: a partially prefilled sequence is always first, because only waiting[0] may be chunked.For each sequence, the new tokens are [num_cached_tokens, num_cached_tokens + num_scheduled_tokens). The function builds:
- input_ids, positions
- The new tokens of all sequences, concatenated, and their absolute positions. Positions start at
start, so a second chunk or a prefix-cache hit continues the numbering RoPE needs. - cu_seqlens_q
- Running sum of new-token counts: where each sequence's queries begin in the packed batch.
- cu_seqlens_k
- Running sum of
seqlen_k = end: each sequence attends to its cached prefix plus its new tokens. - max_seqlen_q, max_seqlen_k
- Longest of each; FlashAttention needs them to size its work.
- slot_mapping
- One slot per new token. The loop visits the blocks the range touches; the first block starts at
start % block_size, the last ends atend − i × block_size, and blocks in between contribute all their slots. - block_tables
- Only when
cu_seqlens_k[-1] > cu_seqlens_q[-1], meaning at least one sequence has cached context. Then attention must read K/V for those earlier positions from the cache, so every sequence's table is passed. OtherwiseNone, and attention uses the freshly computedkandv.
Each list becomes a pinned CPU tensor copied to the GPU with non_blocking=True, so the copies overlap with Python work. set_context receives everything positionally; context_lens is None in prefill.
129 def prepare_prefill(self, seqs: list[Sequence]):130 input_ids = []131 positions = []132 cu_seqlens_q = [0]133 cu_seqlens_k = [0]134 max_seqlen_q = 0135 max_seqlen_k = 0136 slot_mapping = []137 block_tables = None138 for seq in seqs:139 start = seq.num_cached_tokens140 seqlen_q = seq.num_scheduled_tokens141 end = start + seqlen_q142 seqlen_k = end143 input_ids.extend(seq[start:end])144 positions.extend(range(start, end))145 cu_seqlens_q.append(cu_seqlens_q[-1] + seqlen_q)146 cu_seqlens_k.append(cu_seqlens_k[-1] + seqlen_k)147 max_seqlen_q = max(seqlen_q, max_seqlen_q)148 max_seqlen_k = max(seqlen_k, max_seqlen_k)149 if not seq.block_table: # warmup150 continue151 start_block = start // self.block_size152 end_block = (end + self.block_size - 1) // self.block_size153 for i in range(start_block, end_block):154 slot_start = seq.block_table[i] * self.block_size155 if i == start_block:156 slot_start += start % self.block_size157 if i != end_block - 1:158 slot_end = seq.block_table[i] * self.block_size + self.block_size159 else:160 slot_end = seq.block_table[i] * self.block_size + end - i * self.block_size161 slot_mapping.extend(range(slot_start, slot_end))162 if cu_seqlens_k[-1] > cu_seqlens_q[-1]: # prefix cache163 block_tables = self.prepare_block_tables(seqs)164 input_ids = torch.tensor(input_ids, dtype=torch.int64, pin_memory=True).cuda(non_blocking=True)165 positions = torch.tensor(positions, dtype=torch.int64, pin_memory=True).cuda(non_blocking=True)166 cu_seqlens_q = torch.tensor(cu_seqlens_q, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)167 cu_seqlens_k = torch.tensor(cu_seqlens_k, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)168 slot_mapping = torch.tensor(slot_mapping, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)169 set_context(True, cu_seqlens_q, cu_seqlens_k, max_seqlen_q, max_seqlen_k, slot_mapping, None, block_tables)170 return input_ids, positions7.6prepare_decode: one token per sequence
input_idsis each sequence'slast_token, the token sampled in the previous step.positionsislen(seq) − 1, the position that token occupies.context_lensislen(seq). It includes the new token, whose K/V the store kernel writes in this same forward before attention reads the cache.slot_mappingis the last block's base slot pluslast_block_num_tokens − 1.- Block tables are always passed; decode always reads K/V from the cache.
172 def prepare_decode(self, seqs: list[Sequence]):173 input_ids = []174 positions = []175 slot_mapping = []176 context_lens = []177 for seq in seqs:178 input_ids.append(seq.last_token)179 positions.append(len(seq) - 1)180 context_lens.append(len(seq))181 slot_mapping.append(seq.block_table[-1] * self.block_size + seq.last_block_num_tokens - 1)182 input_ids = torch.tensor(input_ids, dtype=torch.int64, pin_memory=True).cuda(non_blocking=True)183 positions = torch.tensor(positions, dtype=torch.int64, pin_memory=True).cuda(non_blocking=True)184 slot_mapping = torch.tensor(slot_mapping, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)185 context_lens = torch.tensor(context_lens, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)186 block_tables = self.prepare_block_tables(seqs)187 set_context(False, slot_mapping=slot_mapping, context_lens=context_lens, block_tables=block_tables)188 return input_ids, positions7.7prepare_block_tables and prepare_sample
Block tables are padded with −1 to the longest table in the batch and sent as int32. The kernels only read as many entries as each sequence's length requires, so the padding is never dereferenced. Temperatures become a float32 tensor, built only on rank 0, the only rank that samples.
123 def prepare_block_tables(self, seqs: list[Sequence]):124 max_len = max(len(seq.block_table) for seq in seqs)125 block_tables = [seq.block_table + [-1] * (max_len - len(seq.block_table)) for seq in seqs]126 block_tables = torch.tensor(block_tables, dtype=torch.int32, pin_memory=True).cuda(non_blocking=True)127 return block_tables190 def prepare_sample(self, seqs: list[Sequence]):191 temperatures = [seq.temperature for seq in seqs]192 temperatures = torch.tensor(temperatures, dtype=torch.float32, pin_memory=True).cuda(non_blocking=True)193 return temperatures7.8run and run_model: eager or graph replay
run is the method rank 0 broadcasts each step: prepare the inputs, compute logits, sample on rank 0, reset the context. Workers return None.
run_model chooses how to execute the model:
- Eager for every prefill, whenever
enforce_eageris set, and for decode batches larger than 512. - Graph replay otherwise. Pick the smallest captured batch size that is at least
bs, copy the real inputs into the firstbsrows of the static buffers, and replay.slot_mappingis filled with −1 first, so padded rows store nothing (the Triton kernel returns early on −1), andcontext_lensis zeroed, so padded rows attend to nothing.block_tablesis not cleared; stale entries are harmless becausecontext_lensbounds what is read. - The LM head runs after the replay on only the real
bsrows, so padding never pays for the 151,936-wide vocabulary projection.
195 @torch.inference_mode()196 def run_model(self, input_ids: torch.Tensor, positions: torch.Tensor, is_prefill: bool):197 if is_prefill or self.enforce_eager or input_ids.size(0) > 512:198 return self.model.compute_logits(self.model(input_ids, positions))199 else:200 bs = input_ids.size(0)201 context = get_context()202 graph = self.graphs[next(x for x in self.graph_bs if x >= bs)]203 graph_vars = self.graph_vars204 graph_vars["input_ids"][:bs] = input_ids205 graph_vars["positions"][:bs] = positions206 graph_vars["slot_mapping"].fill_(-1)207 graph_vars["slot_mapping"][:bs] = context.slot_mapping208 graph_vars["context_lens"].zero_()209 graph_vars["context_lens"][:bs] = context.context_lens210 graph_vars["block_tables"][:bs, :context.block_tables.size(1)] = context.block_tables211 graph.replay()212 return self.model.compute_logits(graph_vars["outputs"][:bs])213 214 def run(self, seqs: list[Sequence], is_prefill: bool) -> list[int]:215 input_ids, positions = self.prepare_prefill(seqs) if is_prefill else self.prepare_decode(seqs)216 temperatures = self.prepare_sample(seqs) if self.rank == 0 else None217 logits = self.run_model(input_ids, positions, is_prefill)218 token_ids = self.sampler(logits, temperatures).tolist() if self.rank == 0 else None219 reset_context()220 return token_ids7.9capture_cudagraph: one graph per batch bucket
A CUDA graph records a fixed sequence of kernel launches with fixed tensor addresses and replays it with a single launch, removing per-kernel CPU overhead. Decode is where that overhead dominates: a decode forward is hundreds of small kernels for a few tokens.
- Static buffers sized for the largest batch:
input_ids,positions,slot_mapping,context_lens,block_tables(widthceil(max_model_len / block_size), 16 by default) andoutputs([max_bs, hidden_size]). - Buckets: 1, 2, 4, 8, then every multiple of 16 up to
max_bs = min(max_num_seqs, 512). That is 36 graphs with the defaults. - Largest first. All graphs share one memory pool, taken from the first (largest) capture, so smaller graphs reuse its memory instead of each reserving their own.
- For each size: point the context at slices of the static buffers, run once eagerly as a warmup, then capture the same call. Because
Attentionand the LM head read the context at Python level during capture, the recorded kernels are bound to those buffer addresses, and replay picks up whatever the buffers hold at that moment.
222 @torch.inference_mode()223 def capture_cudagraph(self):224 config = self.config225 hf_config = config.hf_config226 max_bs = min(self.config.max_num_seqs, 512)227 max_num_blocks = (config.max_model_len + self.block_size - 1) // self.block_size228 input_ids = torch.zeros(max_bs, dtype=torch.int64)229 positions = torch.zeros(max_bs, dtype=torch.int64)230 slot_mapping = torch.zeros(max_bs, dtype=torch.int32)231 context_lens = torch.zeros(max_bs, dtype=torch.int32)232 block_tables = torch.zeros(max_bs, max_num_blocks, dtype=torch.int32)233 outputs = torch.zeros(max_bs, hf_config.hidden_size)234 self.graph_bs = [1, 2, 4, 8] + list(range(16, max_bs + 1, 16))235 self.graphs = {}236 self.graph_pool = None237 238 for bs in reversed(self.graph_bs):239 graph = torch.cuda.CUDAGraph()240 set_context(False, slot_mapping=slot_mapping[:bs], context_lens=context_lens[:bs], block_tables=block_tables[:bs])241 outputs[:bs] = self.model(input_ids[:bs], positions[:bs]) # warmup242 with torch.cuda.graph(graph, self.graph_pool):243 outputs[:bs] = self.model(input_ids[:bs], positions[:bs]) # capture244 if self.graph_pool is None:245 self.graph_pool = graph.pool()246 self.graphs[bs] = graph247 torch.cuda.synchronize()248 reset_context()249 250 self.graph_vars = dict(251 input_ids=input_ids,252 positions=positions,253 slot_mapping=slot_mapping,254 context_lens=context_lens,255 block_tables=block_tables,256 outputs=outputs,257 )Try max_num_seqs = 100 and a batch of 97–100: no bucket is large enough and the lookup raises StopIteration (Step 14).
7.10exit
Close the shared-memory handle, meet at a barrier so every rank has closed it, and let rank 0 unlink it. Drop the captured graphs and their memory pool (added in commit ffafaeb, “Release CUDA Graphs resource before exit”), synchronize, and destroy the process group.
50 def exit(self):51 if self.world_size > 1:52 self.shm.close()53 dist.barrier()54 if self.rank == 0:55 self.shm.unlink()56 if not self.enforce_eager:57 del self.graphs, self.graph_pool58 torch.cuda.synchronize()59 dist.destroy_process_group()Step 08The side channel: Context
The model's forward takes only input_ids and positions, like a Hugging Face model. Everything else attention needs travels through one module-level global.
set_context replaces the global with a new Context; get_context returns it; reset_context puts back an empty one after every run. Keeping the metadata out of the call signature also suits CUDA graphs: during capture, the context holds slices of the static buffers, and the recorded kernels read those buffers on every replay.
1from dataclasses import dataclass2import torch3 4 5@dataclass(slots=True)6class Context:7 is_prefill: bool = False8 cu_seqlens_q: torch.Tensor | None = None9 cu_seqlens_k: torch.Tensor | None = None10 max_seqlen_q: int = 011 max_seqlen_k: int = 012 slot_mapping: torch.Tensor | None = None13 context_lens: torch.Tensor | None = None14 block_tables: torch.Tensor | None = None15 16_CONTEXT = Context()17 18def get_context():19 return _CONTEXT20 21def set_context(is_prefill, cu_seqlens_q=None, cu_seqlens_k=None, max_seqlen_q=0, max_seqlen_k=0, slot_mapping=None, context_lens=None, block_tables=None):22 global _CONTEXT23 _CONTEXT = Context(is_prefill, cu_seqlens_q, cu_seqlens_k, max_seqlen_q, max_seqlen_k, slot_mapping, context_lens, block_tables)24 25def reset_context():26 global _CONTEXT27 _CONTEXT = Context()| Field | Prefill | Decode | Read by |
|---|---|---|---|
is_prefill | True | False | Attention, ParallelLMHead |
cu_seqlens_q | [B+1] int32 | None | attention; LM head picks last rows |
cu_seqlens_k | [B+1] int32 | None | attention |
max_seqlen_q / _k | ints | 0 | attention |
slot_mapping | [N] int32 | [B] int32 | Triton store kernel |
context_lens | None | [B] int32 | flash_attn_with_kvcache |
block_tables | only with cached context | always | both FlashAttention calls |
Writers: prepare_prefill, prepare_decode and capture_cudagraph.
Step 09The model: models/qwen3.py
Qwen3 is a standard pre-norm decoder-only transformer with grouped-query attention, RoPE, SwiGLU MLPs and per-head RMSNorm on queries and keys. This file wires the layer classes together in four modules, from the outside in.
9.1The forward pass at a glance
Qwen3ForCausalLM.forward; the LM head is kept outside it so that the CUDA graph's static output buffer is only [max_bs, hidden_size].Qwen3-0.6B config.json | Value | Where it matters |
|---|---|---|
hidden_size | 1024 | width of the residual stream |
num_hidden_layers | 28 | decoder layers, KV cache layers |
num_attention_heads / num_key_value_heads | 16 / 8 | GQA: two query heads per KV head |
head_dim | 128 | larger than 1024 / 16 = 64, so it must be read from the config |
intermediate_size | 3072 | MLP width |
vocab_size | 151,936 | embedding rows, logits width |
max_position_embeddings | 40,960 | RoPE table length; upper clamp for max_model_len |
rope_theta | 1,000,000 | RoPE base |
attention_bias | false | selects the QK-norm path |
tie_word_embeddings | true | LM head shares the embedding matrix |
rms_norm_eps | 1e-6 | every RMSNorm |
torch_dtype | bfloat16 | default dtype while building the model |
9.2Qwen3ForCausalLM: the outer shell
packed_modules_mappingtells the loader that checkpoint tensors namedq_proj,k_proj,v_projgo into slices of the fusedqkv_proj, andgate_proj,up_projinto slices ofgate_up_proj(Step 11).- With tied embeddings,
lm_head.weight.datais pointed at the embedding's storage. One matrix serves both, so it is stored once. forwardstops at the final norm and returns hidden states;compute_logitsis a separate call. The runner relies on that split: onlyforwardis captured in CUDA graphs.
186class Qwen3ForCausalLM(nn.Module):187 packed_modules_mapping = {188 "q_proj": ("qkv_proj", "q"),189 "k_proj": ("qkv_proj", "k"),190 "v_proj": ("qkv_proj", "v"),191 "gate_proj": ("gate_up_proj", 0),192 "up_proj": ("gate_up_proj", 1),193 }194 195 def __init__(196 self,197 config: Qwen3Config198 ) -> None:199 super().__init__()200 self.model = Qwen3Model(config)201 self.lm_head = ParallelLMHead(config.vocab_size, config.hidden_size)202 if config.tie_word_embeddings:203 self.lm_head.weight.data = self.model.embed_tokens.weight.data204 205 def forward(206 self,207 input_ids: torch.Tensor,208 positions: torch.Tensor,209 ) -> torch.Tensor:210 return self.model(input_ids, positions)211 212 def compute_logits(213 self,214 hidden_states: torch.Tensor,215 ) -> torch.Tensor:216 return self.lm_head(hidden_states)9.3Qwen3Model: embed, 28 layers, norm
The loop threads two tensors through the layers: hidden_states, the latest sublayer output, and residual, the running sum. residual starts as None. The final norm is called with the residual, so it performs the last pending add before normalising; its returned residual is discarded.
162class Qwen3Model(nn.Module):163 164 def __init__(165 self,166 config: Qwen3Config,167 ) -> None:168 super().__init__()169 self.embed_tokens = VocabParallelEmbedding(config.vocab_size, config.hidden_size)170 self.layers = nn.ModuleList([Qwen3DecoderLayer(config) for _ in range(config.num_hidden_layers)])171 self.norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)172 173 def forward(174 self,175 input_ids: torch.Tensor,176 positions: torch.Tensor,177 ) -> torch.Tensor:178 hidden_states = self.embed_tokens(input_ids)179 residual = None180 for layer in self.layers:181 hidden_states, residual = layer(positions, hidden_states, residual)182 hidden_states, _ = self.norm(hidden_states, residual)183 return hidden_states9.4Qwen3DecoderLayer: fused residual adds
x + sublayer(x) explicitly: RMSNorm.add_rms_forward does the add and the norm in one compiled kernel. Layer 0 has no residual yet, so it calls the plain rms_forward and uses the embedding as the residual. r out and h out are the inputs of layer i+1; the last pending add happens in Qwen3Model.norm.- Configuration is read with
getattrdefaults so the same class can load Qwen2-style configs:attention_biasdefaults to True (Qwen2 has QKV bias),head_dimtoNone(then derived),rope_thetato 1,000,000. - Layer 0 is the only one called with
residual=None. It normalises the embedding and keeps the raw embedding as the residual. - Every other norm call is
norm(h, residual), which adds and normalises in one compiled kernel and returns both results (Step 10.4). - The layer returns the MLP output without adding it; the next layer's
input_layernormdoes that.
120class Qwen3DecoderLayer(nn.Module):121 122 def __init__(123 self,124 config: Qwen3Config,125 ) -> None:126 super().__init__()127 self.self_attn = Qwen3Attention(128 hidden_size=config.hidden_size,129 num_heads=config.num_attention_heads,130 num_kv_heads=config.num_key_value_heads,131 max_position=config.max_position_embeddings,132 rms_norm_eps=config.rms_norm_eps,133 qkv_bias=getattr(config, 'attention_bias', True),134 head_dim=getattr(config, 'head_dim', None),135 rope_theta=getattr(config, "rope_theta", 1000000),136 rope_scaling=getattr(config, "rope_scaling", None),137 )138 self.mlp = Qwen3MLP(139 hidden_size=config.hidden_size,140 intermediate_size=config.intermediate_size,141 hidden_act=config.hidden_act,142 )143 self.input_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)144 self.post_attention_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)145 146 def forward(147 self,148 positions: torch.Tensor,149 hidden_states: torch.Tensor,150 residual: torch.Tensor | None,151 ) -> tuple[torch.Tensor, torch.Tensor]:152 if residual is None:153 hidden_states, residual = self.input_layernorm(hidden_states), hidden_states154 else:155 hidden_states, residual = self.input_layernorm(hidden_states, residual)156 hidden_states = self.self_attn(positions, hidden_states)157 hidden_states, residual = self.post_attention_layernorm(hidden_states, residual)158 hidden_states = self.mlp(hidden_states)159 return hidden_states, residual9.5Qwen3Attention: projections, QK-norm, RoPE
Qwen3Attention.forward plus Attention.forward, with Qwen3-0.6B head counts at tp = 1 (16 query heads, 8 KV heads, head_dim 128). The store always happens first, so the attention kernel can treat the new tokens' keys as part of the cache. The dashed path is used for every decode step and for prefill steps where some sequence has cached context.- Head counts are divided by
tp_size; both must divide evenly.q_sizeandkv_sizeare per-rank widths used to split the fused projection output. scaling = head_dim ** -0.5is the softmax scale passed to FlashAttention.o_projmapstotal_num_heads × head_dim(2048) back tohidden_size(1024) and is row-parallel, so it ends with the block's singleall_reduce.- If
rope_scalingis a dict, only itsrope_thetais read. No RoPE scaling method (YaRN, linear, dynamic) is implemented. get_ropeis memoised, so all 28 layers share oneRotaryEmbeddingand one cos/sin table.q_normandk_normexist only when there is no QKV bias. That is Qwen3's QK-norm: an RMSNorm over each head's 128 values, applied before RoPE.- In
forward, the 2-D projection output is viewed as[N, heads, head_dim].o.flatten(1, -1)turns the attention output back into rows; it handles both the prefill shape[N, H, D]and the decode shape[B, 1, H, D].
14class Qwen3Attention(nn.Module):15 16 def __init__(17 self,18 hidden_size: int,19 num_heads: int,20 num_kv_heads: int,21 max_position: int = 4096 * 32,22 head_dim: int | None = None,23 rms_norm_eps: float = 1e-06,24 qkv_bias: bool = False,25 rope_theta: float = 10000,26 rope_scaling: dict | None = None,27 ) -> None:28 super().__init__()29 tp_size = dist.get_world_size()30 self.total_num_heads = num_heads31 assert self.total_num_heads % tp_size == 032 self.num_heads = self.total_num_heads // tp_size33 self.total_num_kv_heads = num_kv_heads34 assert self.total_num_kv_heads % tp_size == 035 self.num_kv_heads = self.total_num_kv_heads // tp_size36 self.head_dim = head_dim or hidden_size // self.total_num_heads37 self.q_size = self.num_heads * self.head_dim38 self.kv_size = self.num_kv_heads * self.head_dim39 self.scaling = self.head_dim ** -0.540 self.qkv_bias = qkv_bias41 42 self.qkv_proj = QKVParallelLinear(43 hidden_size,44 self.head_dim,45 self.total_num_heads,46 self.total_num_kv_heads,47 bias=qkv_bias,48 )49 self.o_proj = RowParallelLinear(50 self.total_num_heads * self.head_dim,51 hidden_size,52 bias=False,53 )54 if isinstance(rope_scaling, dict):55 rope_theta = rope_scaling.get("rope_theta", rope_theta)56 self.rotary_emb = get_rope(57 self.head_dim,58 rotary_dim=self.head_dim,59 max_position=max_position,60 base=rope_theta,61 )62 self.attn = Attention(63 self.num_heads,64 self.head_dim,65 self.scaling,66 self.num_kv_heads,67 )68 if not self.qkv_bias:69 self.q_norm = RMSNorm(self.head_dim, eps=rms_norm_eps)70 self.k_norm = RMSNorm(self.head_dim, eps=rms_norm_eps)71 72 def forward(73 self,74 positions: torch.Tensor,75 hidden_states: torch.Tensor,76 ) -> torch.Tensor:77 qkv = self.qkv_proj(hidden_states)78 q, k, v = qkv.split([self.q_size, self.kv_size, self.kv_size], dim=-1)79 q = q.view(-1, self.num_heads, self.head_dim)80 k = k.view(-1, self.num_kv_heads, self.head_dim)81 v = v.view(-1, self.num_kv_heads, self.head_dim)82 if not self.qkv_bias:83 q = self.q_norm(q)84 k = self.k_norm(k)85 q, k = self.rotary_emb(positions, q, k)86 o = self.attn(q, k, v)87 output = self.o_proj(o.flatten(1, -1))88 return output9.6Qwen3MLP: SwiGLU
One fused projection produces [gate | up] side by side (6,144 wide for Qwen3-0.6B at tp = 1), SiluAndMul computes silu(gate) * up, and the row-parallel down_proj returns to the hidden size. Fusing gate and up turns two matmuls into one larger one. Only silu is supported.
91class Qwen3MLP(nn.Module):92 93 def __init__(94 self,95 hidden_size: int,96 intermediate_size: int,97 hidden_act: str,98 ) -> None:99 super().__init__()100 self.gate_up_proj = MergedColumnParallelLinear(101 hidden_size,102 [intermediate_size] * 2,103 bias=False,104 )105 self.down_proj = RowParallelLinear(106 intermediate_size,107 hidden_size,108 bias=False,109 )110 assert hidden_act == "silu"111 self.act_fn = SiluAndMul()112 113 def forward(self, x):114 gate_up = self.gate_up_proj(x)115 x = self.act_fn(gate_up)116 x = self.down_proj(x)117 return x1import torch2from torch import nn3import torch.distributed as dist4from transformers import Qwen3Config5 6from nanovllm.layers.activation import SiluAndMul7from nanovllm.layers.attention import Attention8from nanovllm.layers.layernorm import RMSNorm9from nanovllm.layers.linear import QKVParallelLinear, MergedColumnParallelLinear, RowParallelLinear10from nanovllm.layers.rotary_embedding import get_rope11from nanovllm.layers.embed_head import VocabParallelEmbedding, ParallelLMHeadStep 10The layers: nanovllm/layers/
Seven files of building blocks. Each is small; the ideas they carry are tensor-parallel sharding, the paged-attention interface, and kernel fusion through torch.compile.
10.1linear.py: tensor-parallel linear layers
MergedColumnParallelLinear stores each rank's slice of gate next to its slice of up, SiluAndMul can split its local output in half and still pair the right columns. The attention block follows the same pattern: qkv_proj is column-parallel by heads, attention runs on local heads, o_proj is row-parallel and ends in one all_reduce.All classes store weight as [output_size, input_size], PyTorch's convention for F.linear, which computes x @ W.T. Each attaches its own weight_loader method to its parameters as an attribute. The loader later finds it with getattr(param, "weight_loader") and lets the layer copy just this rank's shard (Step 11).
1import torch2from torch import nn3import torch.nn.functional as F4import torch.distributed as dist5 6 7def divide(numerator, denominator):8 assert numerator % denominator == 09 return numerator // denominator10 11 12class LinearBase(nn.Module):13 14 def __init__(15 self,16 input_size: int,17 output_size: int,18 bias: bool = False,19 tp_dim: int | None = None,20 ):21 super().__init__()22 self.tp_dim = tp_dim23 self.tp_rank = dist.get_rank()24 self.tp_size = dist.get_world_size()25 self.weight = nn.Parameter(torch.empty(output_size, input_size))26 self.weight.weight_loader = self.weight_loader27 if bias:28 self.bias = nn.Parameter(torch.empty(output_size))29 self.bias.weight_loader = self.weight_loader30 else:31 self.register_parameter("bias", None)32 33 def forward(self, x: torch.Tensor) -> torch.Tensor:34 raise NotImplementedError| Class | Shards | Communication | Used for |
|---|---|---|---|
ReplicatedLinear | nothing | none | not used by Qwen3; available for other models |
ColumnParallelLinear | output rows (tp_dim = 0) | none; output stays sharded | base class of the next two |
MergedColumnParallelLinear | each of several stacked outputs separately | none | gate_up_proj |
QKVParallelLinear | query heads and KV heads separately | none | qkv_proj |
RowParallelLinear | input columns (tp_dim = 1) | all_reduce of the output | o_proj, down_proj |
Replicated and column-parallel
A column-parallel layer keeps output_size / tp rows. Its loader narrows the checkpoint tensor to rows [rank × shard, (rank + 1) × shard). Biases are 1-D along the same dimension, so the same narrowing works for them.
37class ReplicatedLinear(LinearBase):38 39 def __init__(40 self,41 input_size: int,42 output_size: int,43 bias: bool = False,44 ):45 super().__init__(input_size, output_size, bias)46 47 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):48 param.data.copy_(loaded_weight)49 50 def forward(self, x: torch.Tensor) -> torch.Tensor:51 return F.linear(x, self.weight, self.bias)52 53 54class ColumnParallelLinear(LinearBase):55 56 def __init__(57 self,58 input_size: int,59 output_size: int,60 bias: bool = False,61 ):62 tp_size = dist.get_world_size()63 super().__init__(input_size, divide(output_size, tp_size), bias, 0)64 65 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):66 param_data = param.data67 shard_size = param_data.size(self.tp_dim)68 start_idx = self.tp_rank * shard_size69 loaded_weight = loaded_weight.narrow(self.tp_dim, start_idx, shard_size)70 param_data.copy_(loaded_weight)71 72 def forward(self, x: torch.Tensor) -> torch.Tensor:73 return F.linear(x, self.weight, self.bias)Merged column-parallel
output_sizes lists the stacked outputs, [intermediate, intermediate] for gate and up. The loader receives each checkpoint tensor with its index (0 for gate, 1 for up), finds where that output's shard lives inside this rank's parameter (sum(output_sizes[:id]) / tp), and copies this rank's chunk of the checkpoint tensor there. Each rank ends up holding [gate shard; up shard], which is why SiluAndMul can split the local output in half.
76class MergedColumnParallelLinear(ColumnParallelLinear):77 78 def __init__(79 self,80 input_size: int,81 output_sizes: list[int],82 bias: bool = False,83 ):84 self.output_sizes = output_sizes85 super().__init__(input_size, sum(output_sizes), bias)86 87 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor, loaded_shard_id: int):88 param_data = param.data89 shard_offset = sum(self.output_sizes[:loaded_shard_id]) // self.tp_size90 shard_size = self.output_sizes[loaded_shard_id] // self.tp_size91 param_data = param_data.narrow(self.tp_dim, shard_offset, shard_size)92 loaded_weight = loaded_weight.chunk(self.tp_size, self.tp_dim)[self.tp_rank]93 param_data.copy_(loaded_weight)QKV-parallel
The output is (heads + 2 × kv_heads) × head_dim in total. Per rank, the parameter is laid out as local query heads, then local key heads, then local value heads. The loader takes a shard id of "q", "k" or "v", computes that region's offset and size, and copies this rank's chunk of the checkpoint tensor into it. Head counts must divide by tp; there is no KV-head replication for tp larger than the KV head count.
96class QKVParallelLinear(ColumnParallelLinear):97 98 def __init__(99 self,100 hidden_size: int,101 head_size: int,102 total_num_heads: int,103 total_num_kv_heads: int | None = None,104 bias: bool = False,105 ):106 tp_size = dist.get_world_size()107 total_num_kv_heads = total_num_kv_heads or total_num_heads108 self.head_size = head_size109 self.num_heads = divide(total_num_heads, tp_size)110 self.num_kv_heads = divide(total_num_kv_heads, tp_size)111 output_size = (total_num_heads + 2 * total_num_kv_heads) * self.head_size112 super().__init__(hidden_size, output_size, bias)113 114 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor, loaded_shard_id: str):115 param_data = param.data116 assert loaded_shard_id in ["q", "k", "v"]117 if loaded_shard_id == "q":118 shard_size = self.num_heads * self.head_size119 shard_offset = 0120 elif loaded_shard_id == "k":121 shard_size = self.num_kv_heads * self.head_size122 shard_offset = self.num_heads * self.head_size123 else:124 shard_size = self.num_kv_heads * self.head_size125 shard_offset = self.num_heads * self.head_size + self.num_kv_heads * self.head_size126 param_data = param_data.narrow(self.tp_dim, shard_offset, shard_size)127 loaded_weight = loaded_weight.chunk(self.tp_size, self.tp_dim)[self.tp_rank]128 param_data.copy_(loaded_weight)Row-parallel
Each rank keeps input_size / tp columns and computes a partial product from its slice of the input, which is exactly the sharded output of the column-parallel layer before it. all_reduce sums the partials. The bias is added on rank 0 only, so the sum contains it once. Its loader copies a 1-D bias whole and narrows 2-D weights along columns; the 1-D branch was added in commit bf99453 to fix a crash when bias=True.
131class RowParallelLinear(LinearBase):132 133 def __init__(134 self,135 input_size: int,136 output_size: int,137 bias: bool = False,138 ):139 tp_size = dist.get_world_size()140 super().__init__(divide(input_size, tp_size), output_size, bias, 1)141 142 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):143 param_data = param.data144 if param_data.ndim == 1:145 param_data.copy_(loaded_weight)146 return147 shard_size = param_data.size(self.tp_dim)148 start_idx = self.tp_rank * shard_size149 loaded_weight = loaded_weight.narrow(self.tp_dim, start_idx, shard_size)150 param_data.copy_(loaded_weight)151 152 def forward(self, x: torch.Tensor) -> torch.Tensor:153 y = F.linear(x, self.weight, self.bias if self.tp_rank == 0 else None)154 if self.tp_size > 1:155 dist.all_reduce(y)156 return y10.2embed_head.py: vocab-parallel embedding and LM head
VocabParallelEmbedding splits the vocabulary rows evenly across ranks (vocab size must divide by tp; 151,936 divides by 1, 2, 4 and 8). With one rank it is a plain F.embedding. With more, each rank remaps in-range ids to local indices, sends out-of-range ids to row 0, zeroes those rows afterwards, and all_reduce sums the results.
1import torch2from torch import nn3import torch.nn.functional as F4import torch.distributed as dist5 6from nanovllm.utils.context import get_context7 8 9class VocabParallelEmbedding(nn.Module):10 11 def __init__(12 self,13 num_embeddings: int,14 embedding_dim: int,15 ):16 super().__init__()17 self.tp_rank = dist.get_rank()18 self.tp_size = dist.get_world_size()19 assert num_embeddings % self.tp_size == 020 self.num_embeddings = num_embeddings21 self.num_embeddings_per_partition = self.num_embeddings // self.tp_size22 self.vocab_start_idx = self.num_embeddings_per_partition * self.tp_rank23 self.vocab_end_idx = self.vocab_start_idx + self.num_embeddings_per_partition24 self.weight = nn.Parameter(torch.empty(self.num_embeddings_per_partition, embedding_dim))25 self.weight.weight_loader = self.weight_loader26 27 def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):28 param_data = param.data29 shard_size = param_data.size(0)30 start_idx = self.tp_rank * shard_size31 loaded_weight = loaded_weight.narrow(0, start_idx, shard_size)32 param_data.copy_(loaded_weight)33 34 def forward(self, x: torch.Tensor):35 if self.tp_size > 1:36 mask = (x >= self.vocab_start_idx) & (x < self.vocab_end_idx)37 x = mask * (x - self.vocab_start_idx)38 y = F.embedding(x, self.weight)39 if self.tp_size > 1:40 y = mask.unsqueeze(1) * y41 dist.all_reduce(y)42 return yParallelLMHead reuses the same sharded weight layout, so a tied checkpoint shares it with the embedding. In prefill, only the last position of each sequence is needed for sampling, so cu_seqlens_q[1:] − 1 selects those rows before the matmul: a prefill of 16,384 tokens across 4 sequences computes 4 logit rows, not 16,384. With several ranks, dist.gather collects every rank's vocabulary slice on rank 0, which concatenates them; other ranks return None.
45class ParallelLMHead(VocabParallelEmbedding):46 47 def __init__(48 self,49 num_embeddings: int,50 embedding_dim: int,51 bias: bool = False,52 ):53 assert not bias54 super().__init__(num_embeddings, embedding_dim)55 56 def forward(self, x: torch.Tensor):57 context = get_context()58 if context.is_prefill:59 last_indices = context.cu_seqlens_q[1:] - 160 x = x[last_indices].contiguous()61 logits = F.linear(x, self.weight)62 if self.tp_size > 1:63 all_logits = [torch.empty_like(logits) for _ in range(self.tp_size)] if self.tp_rank == 0 else None64 dist.gather(logits, all_logits, 0)65 logits = torch.cat(all_logits, -1) if self.tp_rank == 0 else None66 return logits10.3attention.py: storing K/V and calling FlashAttention
The Triton store kernel
One Triton program per token (grid = (N,)). Each program loads its slot from slot_mapping, returns at once if the slot is −1 (CUDA-graph padding rows), and otherwise copies the token's D = num_kv_heads × head_dim key values and value values into row slot of the flattened caches. D is a compile-time constant, and tl.arange(0, D) requires it to be a power of two: 1,024 for Qwen3-0.6B at tp = 1, 512 at tp = 2.
1import torch2from torch import nn3import triton4import triton.language as tl5 6from flash_attn import flash_attn_varlen_func, flash_attn_with_kvcache7from nanovllm.utils.context import get_context8 9 10@triton.jit11def store_kvcache_kernel(12 key_ptr,13 key_stride,14 value_ptr,15 value_stride,16 k_cache_ptr,17 v_cache_ptr,18 slot_mapping_ptr,19 D: tl.constexpr,20):21 idx = tl.program_id(0)22 slot = tl.load(slot_mapping_ptr + idx)23 if slot == -1: return24 key_offsets = idx * key_stride + tl.arange(0, D)25 value_offsets = idx * value_stride + tl.arange(0, D)26 key = tl.load(key_ptr + key_offsets)27 value = tl.load(value_ptr + value_offsets)28 cache_offsets = slot * D + tl.arange(0, D)29 tl.store(k_cache_ptr + cache_offsets, key)30 tl.store(v_cache_ptr + cache_offsets, value)The wrapper checks the memory layout the kernel assumes: the last dimension is contiguous, heads are head_dim apart, and one cache slot is D contiguous values. The row strides of key and value are passed separately because they differ: key comes out of RoPE as a fresh contiguous tensor (row stride D), while value is still a view into the fused QKV output, whose rows are 4,096 wide for Qwen3-0.6B.
33def store_kvcache(key: torch.Tensor, value: torch.Tensor, k_cache: torch.Tensor, v_cache: torch.Tensor, slot_mapping: torch.Tensor):34 N, num_heads, head_dim = key.shape35 D = num_heads * head_dim36 assert key.stride(-1) == 1 and value.stride(-1) == 137 assert key.stride(1) == head_dim and value.stride(1) == head_dim38 assert k_cache.stride(1) == D and v_cache.stride(1) == D39 assert slot_mapping.numel() == N40 store_kvcache_kernel[(N,)](key, key.stride(0), value, value.stride(0), k_cache, v_cache, slot_mapping, D)The Attention module
k_cacheandv_cachestart as empty tensors;allocate_kv_cachereplaces them with views of the pool. While they are empty (during warmup) the store is skipped.- The store runs before attention, so the current tokens' K/V are already in the cache when attention reads it.
- Prefill without cached context:
flash_attn_varlen_funcon the freshq,k,vwithcu_seqlensdescribing the packed sequences. - Prefill with cached context (
block_tablesset): the same function, butkandvare the whole paged caches andblock_tabletells it where each sequence's keys are. Withcausal=Trueand fewer queries than keys, FlashAttention aligns the causal mask to the bottom-right corner, so query i of a sequence sees every key up to its own position. - Decode:
flash_attn_with_kvcachewithqreshaped to[B, 1, heads, head_dim], readingcontext_lenskeys per sequence through the block table. - GQA needs no special handling; FlashAttention accepts fewer KV heads than query heads.
43class Attention(nn.Module):44 45 def __init__(46 self,47 num_heads,48 head_dim,49 scale,50 num_kv_heads,51 ):52 super().__init__()53 self.num_heads = num_heads54 self.head_dim = head_dim55 self.scale = scale56 self.num_kv_heads = num_kv_heads57 self.k_cache = self.v_cache = torch.tensor([])58 59 def forward(self, q: torch.Tensor, k: torch.Tensor, v: torch.Tensor):60 context = get_context()61 k_cache, v_cache = self.k_cache, self.v_cache62 if k_cache.numel() and v_cache.numel():63 store_kvcache(k, v, k_cache, v_cache, context.slot_mapping)64 if context.is_prefill:65 if context.block_tables is not None: # prefix cache66 k, v = k_cache, v_cache67 o = flash_attn_varlen_func(q, k, v,68 max_seqlen_q=context.max_seqlen_q, cu_seqlens_q=context.cu_seqlens_q,69 max_seqlen_k=context.max_seqlen_k, cu_seqlens_k=context.cu_seqlens_k,70 softmax_scale=self.scale, causal=True, block_table=context.block_tables)71 else: # decode72 o = flash_attn_with_kvcache(q.unsqueeze(1), k_cache, v_cache,73 cache_seqlens=context.context_lens, block_table=context.block_tables, 74 softmax_scale=self.scale, causal=True)75 return o10.4layernorm.py: RMSNorm with a fused residual add
RMSNorm scales each vector by the inverse root-mean-square of its elements, then by a learned weight: y = x / sqrt(mean(x²) + eps) × w. There is no mean subtraction and no bias.
- The statistics are computed in float32 and the result is cast back before multiplying by the weight, matching Hugging Face's Qwen3 implementation.
- The in-place
mul_calls work on the float32 copy, never on the caller's tensor. add_rms_forwardadds the residual in float32, keeps the sum, cast to model dtype, as the new residual, and normalises it. Under@torch.compilethe add, the reduction and the scaling fuse into one or two generated kernels instead of a chain of separate ones.- The same class normalises per head for QK-norm: it reduces over the last dimension, whatever the leading shape.
1import torch2from torch import nn3 4 5class RMSNorm(nn.Module):6 7 def __init__(8 self,9 hidden_size: int,10 eps: float = 1e-6,11 ) -> None:12 super().__init__()13 self.eps = eps14 self.weight = nn.Parameter(torch.ones(hidden_size))15 16 @torch.compile17 def rms_forward(18 self,19 x: torch.Tensor,20 ) -> torch.Tensor:21 orig_dtype = x.dtype22 x = x.float()23 var = x.pow(2).mean(dim=-1, keepdim=True)24 x.mul_(torch.rsqrt(var + self.eps))25 x = x.to(orig_dtype).mul_(self.weight)26 return x27 28 @torch.compile29 def add_rms_forward(30 self,31 x: torch.Tensor,32 residual: torch.Tensor,33 ) -> tuple[torch.Tensor, torch.Tensor]:34 orig_dtype = x.dtype35 x = x.float().add_(residual.float())36 residual = x.to(orig_dtype)37 var = x.pow(2).mean(dim=-1, keepdim=True)38 x.mul_(torch.rsqrt(var + self.eps))39 x = x.to(orig_dtype).mul_(self.weight)40 return x, residual41 42 def forward(43 self,44 x: torch.Tensor,45 residual: torch.Tensor | None = None,46 ) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]:47 if residual is None:48 return self.rms_forward(x)49 else:50 return self.add_rms_forward(x, residual)10.5rotary_embedding.py: RoPE
Rotary embeddings rotate pairs of query and key components by an angle proportional to the token's position, so the dot product between a query and a key depends on their distance. Frequencies are θ_j = base^(−2j / d) for j = 0 … d/2 − 1.
- The constructor precomputes
cosandsinfor every position up tomax_position_embeddingsand stores them concatenated in a[max_pos, 1, head_dim]buffer (cos in the first half, sin in the second). The size-1 axis broadcasts over heads.persistent=Falsekeeps it out of the state dict. forwardgathers rows bypositionsand rotatesqandk. Positions come from the runner, so chunked prefill and prefix hits get the right angles with no extra work.apply_rotary_embpairs elementjwith elementj + d/2(the rotate-half layout used by GPT-NeoX and Hugging Face Qwen), not adjacent elements. It computes in float32.rotary_dimmust equalhead_size; partial rotary is not supported.@lru_cache(1)onget_ropereturns the same module for identical arguments, so the table is built once per model.
1from functools import lru_cache2import torch3from torch import nn4 5 6def apply_rotary_emb(7 x: torch.Tensor,8 cos: torch.Tensor,9 sin: torch.Tensor,10) -> torch.Tensor:11 x1, x2 = torch.chunk(x.float(), 2, dim=-1)12 y1 = x1 * cos - x2 * sin13 y2 = x2 * cos + x1 * sin14 return torch.cat((y1, y2), dim=-1).to(x.dtype)15 16 17class RotaryEmbedding(nn.Module):18 19 def __init__(20 self,21 head_size: int,22 rotary_dim: int,23 max_position_embeddings: int,24 base: float,25 ) -> None:26 super().__init__()27 self.head_size = head_size28 assert rotary_dim == head_size29 inv_freq = 1.0 / (base**(torch.arange(0, rotary_dim, 2, dtype=torch.float) / rotary_dim))30 t = torch.arange(max_position_embeddings, dtype=torch.float)31 freqs = torch.einsum("i,j -> ij", t, inv_freq)32 cos = freqs.cos()33 sin = freqs.sin()34 cache = torch.cat((cos, sin), dim=-1).unsqueeze_(1)35 self.register_buffer("cos_sin_cache", cache, persistent=False)36 37 @torch.compile38 def forward(39 self,40 positions: torch.Tensor,41 query: torch.Tensor,42 key: torch.Tensor,43 ) -> tuple[torch.Tensor, torch.Tensor]:44 cos_sin = self.cos_sin_cache[positions]45 cos, sin = cos_sin.chunk(2, dim=-1)46 query = apply_rotary_emb(query, cos, sin)47 key = apply_rotary_emb(key, cos, sin)48 return query, key49 50 51@lru_cache(1)52def get_rope(53 head_size: int,54 rotary_dim: int,55 max_position: int,56 base: float,57):58 rotary_emb = RotaryEmbedding(head_size, rotary_dim, max_position, base)59 return rotary_emb10.6activation.py: SiluAndMul
Splits the last dimension in half and returns silu(first) × second. Compiled, the split, the activation and the product become a single elementwise kernel.
1import torch2from torch import nn3import torch.nn.functional as F4 5 6class SiluAndMul(nn.Module):7 8 @torch.compile9 def forward(self, x: torch.Tensor) -> torch.Tensor:10 x, y = x.chunk(2, -1)11 return F.silu(x) * y10.7sampler.py: temperature and the exponential race
Logits are divided by each row's temperature and turned into probabilities with softmax. The token is then argmax(p_i / E_i), where each E_i is an independent Exp(1) draw. This picks token i with probability exactly p_i: E_i / p_i is exponential with rate p_i, and the smallest of several independent exponentials is the one with rate p_i with probability p_i / Σp = p_i. It is the Gumbel-max trick written without logarithms.
Everything is elementwise plus one argmax, so the whole sampler compiles into a couple of kernels and avoids torch.multinomial. clamp_min_(1e-10) guards against dividing by a zero draw.
1import torch2from torch import nn3 4 5class Sampler(nn.Module):6 7 @torch.compile8 def forward(self, logits: torch.Tensor, temperatures: torch.Tensor):9 logits = logits.float().div_(temperatures.unsqueeze(dim=1))10 probs = torch.softmax(logits, dim=-1)11 sample_tokens = probs.div_(torch.empty_like(probs).exponential_(1).clamp_min_(1e-10)).argmax(dim=-1)12 return sample_tokensStep 11Weight loading: utils/loader.py
Twenty-eight lines map Hugging Face checkpoint names onto nano-vllm's fused, sharded parameters.
- Read the model's
packed_modules_mapping(empty for models without one). - Open every
*.safetensorsfile in the model directory on the CPU and iterate over its tensor names. - If a mapping key is a substring of the name, rewrite the name to the fused parameter, fetch it, and call its
weight_loaderwith the shard id.breakstops at the first matching key. - The
for … elsebranch runs when no key matched: fetch the parameter by its own name and use itsweight_loader, ordefault_weight_loader(a plain copy) when it has none, as RMSNorm weights do.
1import os2from glob import glob3import torch4from torch import nn5from safetensors import safe_open6 7 8def default_weight_loader(param: nn.Parameter, loaded_weight: torch.Tensor):9 param.data.copy_(loaded_weight)10 11 12def load_model(model: nn.Module, path: str):13 packed_modules_mapping = getattr(model, "packed_modules_mapping", {})14 for file in glob(os.path.join(path, "*.safetensors")):15 with safe_open(file, "pt", "cpu") as f:16 for weight_name in f.keys():17 for k in packed_modules_mapping:18 if k in weight_name:19 v, shard_id = packed_modules_mapping[k]20 param_name = weight_name.replace(k, v)21 param = model.get_parameter(param_name)22 weight_loader = getattr(param, "weight_loader")23 weight_loader(param, f.get_tensor(weight_name), shard_id)24 break25 else:26 param = model.get_parameter(weight_name)27 weight_loader = getattr(param, "weight_loader", default_weight_loader)28 weight_loader(param, f.get_tensor(weight_name))| Checkpoint tensor | Parameter | Loader |
|---|---|---|
model.layers.3.self_attn.q_proj.weight | …qkv_proj.weight, shard "q" | QKVParallelLinear.weight_loader |
…self_attn.k_proj.weight / v_proj | …qkv_proj.weight, shard "k" / "v" | QKVParallelLinear.weight_loader |
…mlp.gate_proj.weight / up_proj | …gate_up_proj.weight, shard 0 / 1 | MergedColumnParallelLinear.weight_loader |
…self_attn.o_proj.weight, …mlp.down_proj.weight | same name | RowParallelLinear.weight_loader |
model.embed_tokens.weight, lm_head.weight | same name | VocabParallelEmbedding.weight_loader |
…q_norm.weight, …input_layernorm.weight, model.norm.weight | same name | default_weight_loader |
model.get_parameterraises if a checkpoint tensor has no matching parameter, so unexpected tensors fail loudly. Parameters missing from the checkpoint are not detected; they keep whatevertorch.emptyleft in memory.- Qwen3-0.6B ships one
model.safetensorsof 1,503,300,328 bytes. That matches about 751.6 M bf16 values: the 596.0 M model parameters plus a separate 155.6 M-valuelm_head.weight, so the tied checkpoint stores the matrix twice. Both names load into the one shared tensor, and whichever comes last is the value used; in a tied checkpoint they are the same matrix. - Tensors are read one at a time through
safe_open(file, "pt", "cpu")and copied straight into the GPU parameters withcopy_.
Step 12One step, end to end, with real shapes
Two concrete steps on Qwen3-0.6B, tp = 1, CUDA graphs on, block size 256. Every row names the code that runs and the shapes it sees.
12.1A prefill step: two new prompts of 300 and 50 tokens
| Stage | Code | What happens |
|---|---|---|
| Schedule | Scheduler.schedule | No cache hits. A needs 2 blocks, B needs 1. 350 tokens fit the 16,384 budget, so both become RUNNING. |
| Inputs | prepare_prefill | input_ids and positions [350] (0…299, 0…49); cu_seqlens_q = cu_seqlens_k = [0, 300, 350]; max_seqlen 300; slot_mapping [350]; block_tables = None. |
| Path | run_model | Prefill is always eager. |
| Embed | VocabParallelEmbedding | [350] → [350, 1024] |
| Per layer × 28 | Qwen3DecoderLayer | qkv_proj [350, 4096] → q [350, 16, 128], k and v [350, 8, 128] → QK-norm, RoPE → store 350 rows of K/V → flash_attn_varlen_func on the fresh tensors → o [350, 16, 128] → o_proj [350, 1024]; MLP [350, 6144] → [350, 3072] → [350, 1024]. |
| Final norm | Qwen3Model.norm | [350, 1024] |
| LM head | ParallelLMHead | Rows cu_seqlens_q[1:] − 1 = [299, 349] → [2, 1024] → logits [2, 151936]. |
| Sample | Sampler | temperatures [2] → token ids [2] → .tolist() |
| Bookkeeping | postprocess | A's block 0 (tokens 0–255) is full and gets hashed; B has no full block. num_cached_tokens becomes 300 and 50; each sequence appends its token (301 and 51 tokens). |
12.2A decode step: three sequences of 300, 1000 and 57 tokens
| Stage | Code | What happens |
|---|---|---|
| Schedule | Scheduler.schedule | Nothing waiting. None of 300, 1000, 57 is ≡ 1 mod 256, so no new blocks. Three sequences, one token each. |
| Inputs | prepare_decode | input_ids [3]; positions = [299, 999, 56]; context_lens = [300, 1000, 57]; slot_mapping [3]; block_tables [3, 4] (tables of 2, 4 and 1 blocks, padded with −1). |
| Path | run_model | bs = 3 → graph for 4. Rows 0–2 get real data; row 3 gets slot −1 and context length 0. |
| Inside the graph | Qwen3ForCausalLM.forward | Same layer sequence with N = 4: the store kernel runs 4 programs and row 3 returns at once; flash_attn_with_kvcache takes q [4, 1, 16, 128] and reads 300, 1000, 57 and 0 keys. Output [4, 1024] lands in graph_vars["outputs"]. |
| LM head | compute_logits(outputs[:3]) | No row selection in decode: [3, 1024] → [3, 151936]. |
| Sample | Sampler | Three token ids. |
| Bookkeeping | postprocess | No block became full (the stored positions were 299, 999 and 56). Each sequence appends a token; any that hit EOS or max_tokens are deallocated and leave running. |
Step 13Where each optimisation lives
The README lists prefix caching, tensor parallelism, torch compilation and CUDA graphs. Here is every performance technique in the code, with its location and the reason it pays.
| Technique | Where | Mechanism | What it saves |
|---|---|---|---|
| Paged KV cache | BlockManager, Attention | Fixed 256-token blocks addressed through block tables | Fragmentation; lets sequences grow without moving data |
| Prefix caching | compute_hash, can_allocate, hash_blocks | Chained xxh64 per full block; shared blocks are ref-counted | Recomputing shared prompt prefixes and memory to hold them twice |
| Lazy eviction | deallocate, _allocate_block | Freed blocks keep their hash until recycled, least recently freed first | Turns finished requests into a cache for future ones |
| Chunked prefill | schedule | The first sequence of a batch may be cut to the token budget | Lets prompts longer than max_num_batched_tokens run at all |
| Continuous batching | Scheduler | The batch is rebuilt every step | Idle GPU time waiting for the slowest request |
| Recompute preemption | preempt | Free a victim's blocks and re-prefill it later, mostly from cache | Deadlock when memory runs out, with no CPU swap space |
| Varlen packing | prepare_prefill | cu_seqlens instead of padding | FLOPs on padding tokens |
| Last-token logits | ParallelLMHead | Select one row per sequence before the vocab matmul | A [tokens × 151,936] matmul in prefill |
| CUDA graphs | capture_cudagraph, run_model | 36 decode graphs, padded to the nearest bucket | CPU launch overhead of hundreds of small kernels per decode step |
torch.compile | RMSNorm, RoPE, SiluAndMul, Sampler | Elementwise chains fused into generated Triton kernels | Memory round-trips between elementwise ops |
| Fused residual + norm | Qwen3DecoderLayer | add_rms_forward does the add and the norm together | A separate read and write of the residual stream per sublayer |
| Fused projections | QKVParallelLinear, MergedColumnParallelLinear | q, k, v in one matmul; gate and up in one matmul | Kernel launches and repeated reads of the input |
| Triton KV store | store_kvcache_kernel | One program per token writes K and V at its slot | Python-level scatter into the paged cache |
| Tensor parallelism | linear.py, embed_head.py | Megatron-style column/row split; two all-reduces per layer | Per-GPU weight memory and per-GPU matmul time |
| Shared-memory IPC | write_shm / read_shm, __getstate__ | One pickle into a shared segment plus Events; decode sends only last tokens | Queue/pipe copies and serialising whole token lists |
| Pinned, async copies | prepare_* | pin_memory=True then .cuda(non_blocking=True) | Blocking host-to-device transfers |
| Measured KV sizing | warmup_model, allocate_kv_cache | Profile the worst prefill, give the rest to the cache | Guesswork that wastes memory or runs out of it |
| Exponential-race sampling | Sampler | argmax(p / Exp(1)) | torch.multinomial |
| Tied embeddings | Qwen3ForCausalLM | LM head points at the embedding's storage | A second 151,936 × 1024 matrix in GPU memory |
Step 14Sharp edges and limits
Behaviour found while reading the code for this guide. The first three are bugs that a particular configuration triggers; the rest are limits worth knowing before you build on the engine.
14.1Request length is not checked
Nothing compares len(prompt) + max_tokens with max_model_len. With CUDA graphs, a decode batch containing a sequence whose block table is wider than ceil(max_model_len / block_size) fails when run_model copies it into the static block_tables buffer. In eager mode the limit is the RoPE table, max_position_embeddings. A sequence that needs more blocks than the entire pool is worse: can_allocate returns −1 forever, and once nothing else is running the decode loop preempts it and assert scheduled_seqs fails.
210 graph_vars["block_tables"][:bs, :context.block_tables.size(1)] = context.block_tables71 assert scheduled_seqs14.2max_num_seqs that is not a multiple of 16 crashes graph lookup
Buckets are 1, 2, 4, 8 and multiples of 16 up to min(max_num_seqs, 512). With max_num_seqs = 100 the largest bucket is 96, but the scheduler can still form decode batches of 97–100. next(x for x in self.graph_bs if x >= bs) then has nothing to yield and raises StopIteration. Values of 8 or less, multiples of 16, or enforce_eager=True avoid it.
202 graph = self.graphs[next(x for x in self.graph_bs if x >= bs)]234 self.graph_bs = [1, 2, 4, 8] + list(range(16, max_bs + 1, 16))14.3Worker ranks ignore a non-default kvcache_block_size
LLMEngine sets Sequence.block_size only in the main process. Spawned workers import sequence.py fresh and keep the class default of 256. On those ranks prepare_decode computes last_block_num_tokens with 256 while the block table was built for the configured size. With kvcache_block_size=512 and a 300-token sequence, rank 0 writes the new token to slot b × 512 + 299 and rank 1 to b × 512 + 43, so ranks ≥ 1 store K/V in the wrong place and outputs go wrong silently. The default block size and tensor_parallel_size=1 are unaffected.
21 Sequence.block_size = config.kvcache_block_size181 slot_mapping.append(seq.block_table[-1] * self.block_size + seq.last_block_num_tokens - 1)14.4Each rank sizes its own KV pool
Every rank runs allocate_kv_cache against its own GPU's free memory, but the scheduler allocates block ids using rank 0's count. If another rank had less free memory (another process on that GPU, for example) it has fewer blocks than rank 0 hands out. On identical, otherwise idle GPUs the counts match.
14.5One tensor-parallel engine per machine
The NCCL rendezvous port (2333) and the shared-memory name ("nanovllm") are hard-coded. A second engine with tp > 1 on the same machine collides with the first. A run that dies without reaching exit can leave the "nanovllm" segment behind, and the next SharedMemory(create=True) then fails with FileExistsError. The segment is 1 MiB; a pickled step must fit in it.
14.6Scheduling policy
- Prefill always wins. While a waiting request can be admitted, no sequence decodes, and a long chunked prompt pauses decoding for every step it takes.
- Strict first-come-first-served admission: a head request that does not fit blocks smaller ones behind it.
runningcan exceedmax_num_seqs; only its firstmax_num_seqssequences decode until some finish.- Preemption always recomputes; there is no swap to CPU memory.
14.7API and model coverage
- Unknown keyword arguments to
LLM(...)are dropped silently (2.1). - No greedy decoding (temperature must be > 1e-10), no top-k, top-p or penalties, no stop strings, no streaming. Output text includes the EOS token.
generatereturns dicts despite itslist[str]annotation. - One architecture: Qwen3, plus Qwen2-style configs with QKV bias. RoPE scaling is not implemented. Head counts and the vocabulary must divide by tp, and
(num_key_value_heads / tp) × head_dimmust be a power of two for the Triton store kernel. - Offline only: one process, one Python call at a time, all ranks on one machine.
- The code reads
hf_config.dtype, which needs a transformers release that exposesdtypeon configs.
Step 15Glossary and next exercises
15.1Glossary
- block
- 256 token slots of K/V for every layer; the unit of allocation.
- block table
- A sequence's list of physical block ids, in logical order.
- slot
- Index of one token's row in the flattened cache:
block_id × block_size + offset. - slot_mapping
- The slot for every token computed in a step; tells the store kernel where to write.
- prefix cache
hash_to_block_id: full blocks indexed by the chained hash of everything up to them.- ref_count
- Number of live sequences whose block table contains a block.
- prefill / decode
- Computing a prompt (many tokens per sequence) versus generating (one token per sequence).
- chunked prefill
- Splitting a prompt across steps when it exceeds the token budget.
- preemption
- Taking a running sequence's blocks away and re-queuing it for recomputation.
- cu_seqlens
- Cumulative sequence lengths that describe a packed varlen batch.
- GQA
- Grouped-query attention: fewer K/V heads than query heads (8 and 16 here).
- QK-norm
- RMSNorm applied to each query and key head before RoPE (Qwen3).
- RoPE
- Rotary position embedding: rotate q and k component pairs by position-dependent angles.
- rank
- One process / GPU in tensor parallelism; rank 0 also schedules and samples.
- column / row parallel
- Splitting a weight matrix by output rows (no communication) or input columns (all-reduce).
- CUDA graph bucket
- A captured batch size; a decode batch replays the smallest bucket that fits it.
- Context
- The module-level object that carries per-step attention metadata into the model.
15.2Exercises that touch every layer of the stack
- Watch the prefix cache. In
Scheduler.schedule, lognum_cached_blocksfor each admitted sequence, then runexample.pywith two prompts that share a long system prompt. - Add greedy decoding. Relax the assertion in
SamplingParams, and inSampler.forwarduseargmaxon rows whose temperature is 0 before dividing. - Add top-k. Carry a
top_kfield fromSamplingParamsthroughSequenceandprepare_sample, and mask logits below each row's k-th largest value before softmax. - Validate request length. Reject or truncate in
add_requestwhenlen(prompt) + max_tokens > max_model_len(14.1). - Fix the bucket lookup. Append
max_bstograph_bswhen it is not already the last bucket (14.2). - Fix worker block size. Set
Sequence.block_size = config.kvcache_block_sizeinModelRunner.__init__(14.3).