Nano-vLLM Field Guide

Code walkthrough · mohan-zeyu/nano-vllm · commit e86692c

Nano-vLLM, line by line

Nano-vLLM serves Qwen3 models with a paged KV cache, prefix caching, chunked prefill, tensor parallelism, CUDA graphs and torch.compile in under 1,400 lines of Python. This guide follows a request from llm.generate() down to the Triton kernel that writes its keys and values, and puts the exact source under every point.

  • 21 Python files
  • 16 steps
  • every line of code shown below its explanation
Contents
  1. Step 00Orientation
    1. 0.1Repository map
    2. 0.2How the pieces stack
    3. 0.3Five ideas the code assumes
    4. 0.4The route through this guide
  2. Step 01The front door
    1. 1.1The public surface: LLM and SamplingParams
    2. 1.2example.py: the smallest real use
    3. 1.3bench.py: the throughput benchmark
    4. 1.4Packaging and dependencies
  3. Step 02Configuration
    1. 2.1Config: every engine knob
    2. 2.2SamplingParams: per-request decoding
  4. Step 03The engine loop: LLMEngine
    1. 3.1Construction, and why the order matters
    2. 3.2add_request: text or token ids in, a Sequence out
    3. 3.3step: one iteration of the serving loop
    4. 3.4generate: batch in, ordered results out
    5. 3.5exit: shutting down the ranks
    6. 3.6The life of one request
  5. Step 04The unit of work: Sequence
    1. 4.1Identity, counters and block table
    2. 4.2Block arithmetic and list behaviour
    3. 4.3Status transitions
    4. 4.4Compact pickling for worker ranks
  6. Step 05Paged KV memory: BlockManager
    1. 5.1Blocks, block tables and slots
    2. 5.2Block and the four bookkeeping structures
    3. 5.3Chained block hashes
    4. 5.4can_allocate: count cache hits, check capacity
    5. 5.5allocate: build the block table
    6. 5.6_allocate_block: take a block, evict its old identity
    7. 5.7deallocate: release a sequence's blocks
    8. 5.8can_append and may_append: growing during decode
    9. 5.9hash_blocks: publish newly full blocks
  7. Step 06Choosing each batch: Scheduler
    1. 6.1Queues and limits
    2. 6.2Prefill phase and chunked prefill
    3. 6.3Decode phase and preemption
    4. 6.4preempt: give memory back, recompute later
    5. 6.5postprocess: apply the sampled tokens
    6. 6.6Simulator: run the scheduler yourself
  8. Step 07From sequences to tensors: ModelRunner
    1. 7.1Boot sequence
    2. 7.2The control plane: call, write_shm, read_shm, loop
    3. 7.3warmup_model: measure the worst prefill
    4. 7.4allocate_kv_cache: size the pool from what is left
    5. 7.5prepare_prefill: pack a varlen batch
    6. 7.6prepare_decode: one token per sequence
    7. 7.7prepare_block_tables and prepare_sample
    8. 7.8run and run_model: eager or graph replay
    9. 7.9capture_cudagraph: one graph per batch bucket
    10. 7.10exit
  9. Step 08The side channel: Context
  10. Step 09The model: models/qwen3.py
    1. 9.1The forward pass at a glance
    2. 9.2Qwen3ForCausalLM: the outer shell
    3. 9.3Qwen3Model: embed, 28 layers, norm
    4. 9.4Qwen3DecoderLayer: fused residual adds
    5. 9.5Qwen3Attention: projections, QK-norm, RoPE
    6. 9.6Qwen3MLP: SwiGLU
  11. Step 10The layers: nanovllm/layers/
    1. 10.1linear.py: tensor-parallel linear layers
    2. 10.2embed_head.py: vocab-parallel embedding and LM head
    3. 10.3attention.py: storing K/V and calling FlashAttention
    4. 10.4layernorm.py: RMSNorm with a fused residual add
    5. 10.5rotary_embedding.py: RoPE
    6. 10.6activation.py: SiluAndMul
    7. 10.7sampler.py: temperature and the exponential race
  12. Step 11Weight loading: utils/loader.py
  13. Step 12One step, end to end, with real shapes
    1. 12.1A prefill step: two new prompts of 300 and 50 tokens
    2. 12.2A decode step: three sequences of 300, 1000 and 57 tokens
  14. Step 13Where each optimisation lives
  15. Step 14Sharp edges and limits
    1. 14.1Request length is not checked
    2. 14.2max_num_seqs that is not a multiple of 16 crashes graph lookup
    3. 14.3Worker ranks ignore a non-default kvcache_block_size
    4. 14.4Each rank sizes its own KV pool
    5. 14.5One tensor-parallel engine per machine
    6. 14.6Scheduling policy
    7. 14.7API and model coverage
  16. Step 15Glossary and next exercises
    1. 15.1Glossary
    2. 15.2Exercises that touch every layer of the stack

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.

FileLinesRole
README.md66Install, model download, quick start, benchmark table
pyproject.toml27Package metadata and the five runtime dependencies
example.py33Minimal chat-style usage
bench.py32Throughput benchmark: 256 random requests
nanovllm/__init__.py2Exports LLM and SamplingParams
nanovllm/llm.py5LLM is an empty subclass of LLMEngine
nanovllm/config.py25Every engine knob, validated in __post_init__
nanovllm/sampling_params.py11temperature, max_tokens, ignore_eos
nanovllm/engine/llm_engine.py90Builds everything; generate() → step() loop
nanovllm/engine/sequence.py83One request: tokens, counters, block table
nanovllm/engine/block_manager.py120Paged KV allocator with prefix caching
nanovllm/engine/scheduler.py92Chooses each step's batch; preemption; bookkeeping
nanovllm/engine/model_runner.py257Per-GPU worker: KV pool, tensors, CUDA graphs, IPC
nanovllm/utils/context.py27Global per-step metadata read inside the model
nanovllm/utils/loader.py28safetensors → sharded parameters
nanovllm/models/qwen3.py216Qwen3 architecture assembled from layers
nanovllm/layers/linear.py156Replicated / column / row / merged / QKV parallel linears
nanovllm/layers/embed_head.py66Vocab-parallel embedding and LM head
nanovllm/layers/attention.py75Triton KV-store kernel + FlashAttention calls
nanovllm/layers/layernorm.py50RMSNorm with fused residual add
nanovllm/layers/rotary_embedding.py59RoPE cos/sin cache
nanovllm/layers/activation.py11SiluAndMul
nanovllm/layers/sampler.py12Temperature + 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

YOUR SCRIPTPUBLIC API · nanovllm/__init__.py · llm.py · sampling_params.pyENGINE · CPU · MAIN PROCESS · nanovllm/engine/RUNNER · ONE PROCESS PER GPU · engine/model_runner.py · utils/MODEL · nanovllm/models/qwen3.pyLAYERS · nanovllm/layers/*.pyexample.py2 chat promptsbench.py256 random requestsLLMclass LLM(LLMEngine): passSamplingParamstemperature · max_tokens · ignore_eosLLMEnginellm_engine.pySchedulerscheduler.pyBlockManagerblock_manager.pySequencesequence.pyModelRunner · rank 0model_runner.py · owns SamplerModelRunner · rank 1…N−1spawned only when tp > 1Contextper-step metadataload_modelutils/loader.pyQwen3ForCausalLMforward() · compute_logits()Qwen3Modelembed → layers → normQwen3DecoderLayer × 28attention block + MLP blocklinear.pyTP matmulsembed_head.pyvocab-parallelattention.pyKV store + attnlayernorm.pyRMSNormrotary_embedding.pyRoPEactivation.pySiluAndMulsampler.pysamplinggenerate(prompts, sampling_params)is ancall('run', seqs, is_prefill) → token_idsmodel(input_ids, positions) · compute_logits()built fromschedule()allocate()block_tableshm + Eventself.modellayers[i]
The call stack, top to bottom. Everything above the runner band is plain Python on the CPU of the main process. The runner band is replicated once per GPU: rank 0 lives inside the main process, ranks 1…N−1 are spawned processes that mirror every call rank 0 makes. 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

  1. Steps 01–03 cover the surface: how the package is used, every configuration knob, and the engine loop that drives everything.
  2. Steps 04–06 are the CPU-side bookkeeping: the Sequence object, 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.
  3. Steps 07–08 cross to the GPU: the per-rank ModelRunner that sizes the KV cache, builds the batch tensors and replays CUDA graphs, and the Context side channel it fills.
  4. Steps 09–11 are the model: Qwen3's module tree, each layer file, and the weight loader that fills sharded parameters.
  5. 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).

nanovllm/__init__.pyL1–2GitHub ↗
1from nanovllm.llm import LLM2from nanovllm.sampling_params import SamplingParams
nanovllm/llm.pyL1–5GitHub ↗
1from nanovllm.engine.llm_engine import LLMEngine2 3 4class LLM(LLMEngine):5    pass

1.2example.py: the smallest real use

  • The model path must be a local directory. Config asserts os.path.isdir(model), so a Hugging Face hub id such as "Qwen/Qwen3-0.6B" will not work. The README's huggingface-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 inside add_request.
  • enforce_eager=True skips CUDA graph capture. Start-up is faster; decode steps are slower (Step 07.9).
  • One SamplingParams object is shared by both prompts; generate broadcasts it.
  • Each output is a dict with "text" and "token_ids", in the same order as the prompts.
example.pyL1–33GitHub ↗
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], and max_tokens drawn from 100–1024. seed(0) makes the workload repeatable.
  • ignore_eos=True makes every request produce exactly max_tokens tokens, so the total output token count is known in advance.
  • enforce_eager=False turns CUDA graphs on; max_model_len=4096 covers 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).
bench.pyL1–32GitHub ↗
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:

DependencyUsed forWhere
torch>=2.4.0tensors, NCCL, CUDA graphs, torch.compile, spawn-based multiprocessingeverywhere
triton>=3.0.0the KV-cache store kernel; also the backend torch.compile generates kernels withlayers/attention.py
transformers>=4.51.0AutoConfig, AutoTokenizer, the Qwen3Config typeconfig.py, llm_engine.py, models/qwen3.py
flash-attnflash_attn_varlen_func (prefill) and flash_attn_with_kvcache (decode)layers/attention.py
xxhash64-bit hashes of KV blocks for prefix cachingengine/block_manager.py

safetensors, numpy and tqdm are imported too; they arrive transitively through transformers.

pyproject.tomlL1–27GitHub ↗
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.

nanovllm/config.pyL1–25GitHub ↗
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)
FieldDefaultRead byEffect
modelrequiredConfig, LLMEngine, loaderLocal checkpoint directory. Source of config.json, the tokenizer and *.safetensors.
max_num_batched_tokens16384Scheduler, warmupToken budget of one prefill step. A prompt longer than this is split into chunks. Also sizes the warmup batch.
max_num_seqs512Scheduler, warmup, CUDA graphsMost sequences in one step, prefill or decode. Largest captured graph is min(max_num_seqs, 512).
max_model_len4096warmup, CUDA graphsClamped 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_utilization0.9allocate_kv_cacheFraction of total GPU memory the process may plan for. Whatever weights, activations and other processes leave inside it becomes KV cache.
tensor_parallel_size1LLMEngine, ModelRunner, all TP layersNumber of GPU processes. Asserted to be 1–8.
enforce_eagerFalseModelRunnerTrue skips CUDA graph capture; every forward runs eagerly.
hf_configNoneModelRunner, modelSet to AutoConfig.from_pretrained(model) in __post_init__.
eos−1SchedulerSet by LLMEngine to tokenizer.eos_token_id.
kvcache_block_size256everything that touches blocksTokens per KV block. Must be a multiple of 256, the page size FlashAttention's paged-KV kernels require.
num_kvcache_blocks−1Scheduler → BlockManagerSet 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.

nanovllm/sampling_params.pyL1–11GitHub ↗
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

MAIN PROCESS · rank 0 · cuda:0WORKER PROCESS · rank i, 1 ≤ i < tp · cuda:iConfig(model, **config_kwargs)AutoConfig.from_pretrained · clamp max_model_lenSequence.block_size = kvcache_block_sizeclass attribute, set in this process onlyspawn ModelRunner(config, i, event)for i in 1 … tp−1 · start method 'spawn'ModelRunner(config, 0, events)init_process_group('nccl', tcp://localhost:2333)build Qwen3ForCausalLM on cuda, hf dtypeload_model(): safetensors → parameterswarmup_model(): measure peak activationsallocate_kv_cache(): size and bind the poolcapture_cudagraph() (skipped if enforce_eager)SharedMemory('nanovllm', create) → barrierAutoTokenizer · config.eos = tokenizer.eos_token_idScheduler(config) → BlockManager(…)ModelRunner(config, i, event)init_process_group(…, rank=i)same build as rank 0:model · weights · warmupits own KV cache · its own CUDA graphsbarrier → SharedMemory('nanovllm') attachloop(): event.wait() → read_shm() → call()__init__ never returns until 'exit' arrivesmp spawnrendezvousbarrierconfig.num_kvcache_blocks
What 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.
  1. Filter kwargs into Config. dataclasses.fields(Config) gives the allowed names; everything else is dropped.
  2. Set Sequence.block_size. It is a class attribute, so every Sequence in this process uses the configured block size for num_blocks and friends. The assignment happens only in the main process (see Step 14 for what that means for workers).
  3. 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 the ModelRunner class itself, so constructing it is the whole life of the worker: its __init__ ends in self.loop(). Each worker gets one Event; rank 0 later receives the list of all of them.
  4. Build rank 0's ModelRunner in this process. It must come after the spawns because init_process_group blocks until all ranks have joined. This call loads the model, measures memory, allocates the KV cache and captures CUDA graphs (Step 07).
  5. Load the tokenizer and record EOS. config.eos is set after the workers were spawned, so their copies of config keep −1. They never need it.
  6. Build the Scheduler last. Its BlockManager needs config.num_kvcache_blocks, which only exists after step 4. Spawned workers received a pickled copy of config earlier, so this value is rank 0's own measurement.
  7. Register exit with atexit, so workers are told to stop when the interpreter exits.
nanovllm/engine/llm_engine.pyL1–35GitHub ↗
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().

nanovllm/engine/llm_engine.pyL43–47GitHub ↗
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

  1. scheduler.schedule() returns the batch and whether it is a prefill batch.
  2. num_tokens is measured before the forward, because postprocess resets every num_scheduled_tokens to 0. The sign carries the phase: positive means “prefill tokens processed”, negative means “number of sequences decoded”. generate uses the sign to pick which throughput figure to update.
  3. model_runner.call("run", seqs, is_prefill) runs the forward on every rank and returns one token id per sequence on rank 0.
  4. scheduler.postprocess appends tokens, registers newly full blocks for prefix caching, and retires finished sequences.
  5. The return value lists (seq_id, completion_token_ids) only for sequences that finished in this step.
nanovllm/engine/llm_engine.pyL49–58GitHub ↗
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 SamplingParams is 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_id and then sorted. seq_id comes from a process-wide itertools.count(), so sorting restores submission order even though requests finish in any order, and it keeps working across repeated generate calls.
  • tokenizer.decode(token_ids) does not skip special tokens. A request that stopped on EOS includes the EOS token in both token_ids and text (for Qwen3 chat models, <|im_end|>).
  • The return annotation says list[str]; the function actually returns a list of dicts.
nanovllm/engine/llm_engine.pyL57–90GitHub ↗
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 outputs

3.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.

nanovllm/engine/llm_engine.pyL37–41GitHub ↗
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

LLMEngine.generateSchedulerBlockManagerModelRunner rank 0Qwen3 + Sampleradd(seq) × len(prompts)loop · while not is_finished(): step()schedule()can_allocate · allocatecan_append · may_append(seqs, is_prefill)call('run', seqs, is_prefill)prepare_*() → set_context()model(input_ids, positions)store K/V → flash-attncompute_logits() → sampler()token_ids (rank 0 only)token_idspostprocess(seqs, token_ids)hash_blocks · deallocatecollect finished (seq_id, completion_token_ids)sort by seq_id → tokenizer.decode → [{'text', 'token_ids'}]
One call to 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

nanovllm/engine/sequence.pyL1–31GitHub ↗
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_token never mutates it.
num_tokens
len(token_ids), kept as a counter because worker processes receive a Sequence without its token list during decode (4.4).
num_prompt_tokens
Fixed at creation. Splits token_ids into 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 by deallocate.
num_scheduled_tokens
How many tokens this step will compute, starting at position num_cached_tokens. Set by the scheduler, read by prepare_* and hash_blocks, reset to 0 in postprocess.
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.
K/V already in cachecomputed this stepgenerated token(a) final prefill chunk · num_tokens = 14012345678910111213block 3block 7block 1block 9num_cached_tokens = 8num_scheduled_tokens = 6start = num_cached_tokens = 8end = 8 + 6 = 14 = num_tokens→ moves to running this step(b) decode step · 14 prompt + 3 generatednum_prompt_tokens = 14completion = 3012345678910111213141516block 3block 7block 1block 9b5num_cached_tokens = 16num_scheduled_tokens = 1input_ids = [last_token]positions = [16]context_lens = [17]slot = 5 × 4 + 1 − 1 = 20
The four counters on a 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 writes len(seq) and seq[start:end].
  • num_blocks is ceil(num_tokens / block_size), written with integer arithmetic.
  • last_block_num_tokens is how many tokens occupy the last block, from 1 to block_size. Decode uses it to find the slot of the newest token.
  • block(i) returns the token ids of logical block i; the last block may be partial. Prefix caching hashes these lists.
  • append_token updates token_ids, last_token and num_tokens together.
nanovllm/engine/sequence.pyL33–70GitHub ↗
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 += 1

4.3Status transitions

nanovllm/engine/sequence.pyL8–11GitHub ↗
8class SequenceStatus(Enum):9    WAITING = auto()10    RUNNING = auto()11    FINISHED = auto()
WAITINGscheduler.waitingRUNNINGscheduler.runningFINISHEDblocks releasedadd_request()schedule()final chunk in batchpostprocess()EOS or max_tokenspreempt(): free blocks, appendleft to waitingpartial chunk: stays at waiting[0]decode step: append_token()
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_prefill slices seq[start:end];
  • only last_token in decode, because prepare_decode needs 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.

nanovllm/engine/sequence.pyL72–83GitHub ↗
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_state

Step 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

LOGICAL VIEW · positions of one sequencePHYSICAL KV POOL · one layer of k_cachelogical block 00123bt[0] = 3logical block 14567bt[1] = 7logical block 2891011bt[2] = 1logical block 31213bt[3] = 90123block 0 · free4567block 1891011block 2 · free12131415block 316171819block 4 · free20212223block 5 · free24252627block 6 · free28293031block 732333435block 8 · free36373839block 9position 13 → bt[13 // 4] = 9 → slot 9 × 4 + 13 % 4 = 37
Paging with a toy block size of 4. The sequence sees contiguous positions; the KV pool stores them in whichever blocks were free. 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.

nanovllm/engine/block_manager.pyL1–33GitHub ↗
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()
StructureTypeMeaning
blockslist[Block]All block objects, indexed by id.
hash_to_block_iddict[int, int]Prefix hash → the block holding that prefix's last block of K/V. This is the prefix cache.
free_block_idsdeque[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_idsset[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.

nanovllm/engine/block_manager.pyL35–41GitHub ↗
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()
seq A7391448265h₀ = H(7 3 9 1)h₁ = H(h₀ ‖ 4 4 8 2)partial: no hashblock 3hash h₀ · ref_count 2block 7hash h₁ · ref_count 2block 1A's tail · ref 1block 9B's tail · ref 1seq B73914482331hithitmiss → new blocksame h₀, same tokenssame h₁ (needs equal h₀ too)A ran first: hash_blocks()registered h₀ and h₁can_allocate(B) → 2num_cached_tokens = 8only 3 tokens computed
Prefix caching keys each full block by a hash chained through every earlier block, so a block only matches when the entire prefix up to and including it matches. B finds h₀ and h₁ in 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.

  1. 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.
  2. 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.
  3. num_new_blocks starts 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.
  4. Return −1 if the free deque is shorter than num_new_blocks, otherwise the number of cached blocks.
nanovllm/engine/block_manager.pyL58–73GitHub ↗
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_blocks

5.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_tokens becomes num_cached_blocks × block_size; prefill will start computing at that position.
nanovllm/engine/block_manager.pyL75–92GitHub ↗
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_size

5.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.

nanovllm/engine/block_manager.pyL43–51GitHub ↗
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_id

5.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.

nanovllm/engine/block_manager.pyL53–56GitHub ↗
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)
nanovllm/engine/block_manager.pyL94–101GitHub ↗
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 timelen % 4Position stored this stepAction (toy block size 4)
807, last slot of block 1none; block 1 becomes full and is hashed in postprocess
918, first slot of block 2may_append allocates block 2
1029none
11310none
12011, block 2 fullnone; block 2 hashed
nanovllm/engine/block_manager.pyL103–108GitHub ↗
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_size tokens.
  • 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.
nanovllm/engine/block_manager.pyL110–120GitHub ↗
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_id

Step 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.

nanovllm/engine/scheduler.pyL1–23GitHub ↗
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.

  1. First come, first served. The loop only looks at waiting[0]. When the head cannot be served it breaks; it never skips ahead to a smaller request.
  2. Budget. remaining is what is left of max_num_batched_tokens. A batch that has used it all stops.
  3. Fresh or continuing. A sequence without blocks is new or was preempted: can_allocate reports 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 after num_cached_tokens.
  4. The chunking rule. If the work does not fit in remaining and 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 than max_num_batched_tokens. With the defaults (16,384 tokens of budget, max_model_len 4,096) that never happens; lower the budget or raise max_model_len to see it.
  5. Commit. Blocks are allocated only now that the sequence is definitely in the batch. num_scheduled_tokens is the work, capped by the budget.
  6. Done or not. If this chunk reaches num_tokens, the sequence becomes RUNNING and moves to running. Otherwise it stays at waiting[0]; it used the whole remaining budget, so the next iteration hits remaining == 0 and the loop ends.
nanovllm/engine/scheduler.pyL25–55GitHub ↗
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, True

6.3Decode phase and preemption

  1. Only reached when nothing could be prefilled.
  2. Take sequences from the front of running, up to max_num_seqs.
  3. Make sure the sequence can store one more token (can_append). While it cannot, preempt the sequence at the back of running, the most recently admitted one, which frees its blocks. If no other sequence is left, preempt this one and move on.
  4. Python's while … else: the else block runs when the loop condition turns false, meaning there is room, and is skipped after the break that follows self-preemption. In the else branch the sequence is set to compute one token, marked is_prefill = False so it pickles compactly, and given a new block if it needs one.
  5. assert scheduled_seqs fires only if a single sequence needs more blocks than the whole pool, because then even preempting everything else frees nothing usable.
  6. running.extendleft(reversed(scheduled_seqs)) puts the batch back at the front of running in its original order.
nanovllm/engine/scheduler.pyL57–73GitHub ↗
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, False

6.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:

  • deallocate keeps hashes, and decode steps hashed every block they filled. When the sequence is re-admitted, can_allocate finds 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 postprocess appends it.
nanovllm/engine/scheduler.pyL75–79GitHub ↗
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

  1. hash_blocks first, while num_cached_tokens and num_scheduled_tokens still describe what was just computed.
  2. Advance num_cached_tokens, clear num_scheduled_tokens.
  3. 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.
  4. Otherwise append the token. Finish when the token is EOS (unless ignore_eos) or the completion has reached max_tokens. The token is appended before the check, so a finishing EOS is part of the output.
  5. A finished sequence frees its blocks immediately and leaves running (an O(n) deque.remove).
nanovllm/engine/scheduler.pyL81–92GitHub ↗
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.

Scheduler and block pool

Block size 4 · toy model · state after each step()

Step
Add
K/V in cachenot computed yetgenerated tokencomputed in the last steppool: in usepool: free but still cached

Sequences

KV block pool

What happened

    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

    1. 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.
    2. 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 every torch.empty inside the layer constructors allocates GPU memory in model precision. Nothing is initialised on the CPU first. Older transformers releases call this attribute torch_dtype; the code switched to dtype in commit 8d63a98, and an install whose config objects lack dtype fails on this line.
    3. Load weights with load_model (Step 11) and create the Sampler.
    4. Warm up before allocating the KV cache, so the measured activation peak does not include it (7.3).
    5. 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.
    6. 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.
    7. 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 enter loop() and stay there.
    nanovllm/engine/model_runner.pyL1–48GitHub ↗
    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

    rank 0 · ModelRunner.call(name, *args)1 write_shm(): pickle [name, *args]2 buf[0:4] = n, little-endian3 buf[4:4+n] = payload4 event.set() for every worker5 run the method locallySharedMemory('nanovllm') · 2²⁰ bytesnpickle payloadunusedwritesrank 1event.wait() → read_shm()→ event.clear() → call()rank 2 … N−1event.wait() → read_shm()→ event.clear() → call()readsEvent.set() · one Event per workerDATA PLANE · NCCL process group via tcp://localhost:2333rank 0 GPUrank 1 GPUrank N−1 GPUall_reducegather → rank 0
    Two channels between ranks. The control plane (top) carries only small pickled messages: which method to run and the 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 through getattr. 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 run until every worker has entered the same run, which they only do after reading and clearing. Rank 0 therefore never overwrites a message a worker has not read yet.
    nanovllm/engine/model_runner.pyL61–89GitHub ↗
    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) and num_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_prefill skips their slot mapping, and Attention skips the KV store because k_cache is still the empty placeholder.
    • The first call of each @torch.compile function 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.
    nanovllm/engine/model_runner.pyL91–101GitHub ↗
    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 − free includes 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.

    nanovllm/engine/model_runner.pyL103–121GitHub ↗
    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 += 1
    kv_cache = torch.empty(2, num_layers, num_blocks, block_size, num_kv_heads, head_dim)[0] keysL0L1L2L3L4L5…L27[1] valuesL0L1L2L3L4L5…L27k_cache of layer 2 = kv_cache[0, 2]block 0block 1block 2block 3per-layer view shape[num_blocks, block_size, kv_heads, head_dim]one cell = one slotkv_heads × head_dim values (D)highlighted slotblock 2 · offset 5 → slot 2 × block_size + 5Triton storek_cache + slot · D + arange(D)binding · a view, no copymodule.k_cache = kv_cache[0, layer_id]
    One tensor holds the whole cache. allocate_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.

    KV pool calculator

    The formula above with your numbers. The two overhead inputs are example values; warmup measures the real ones.

    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

    seq C · 2nd chunkseq A · fresh promptseq B · prefix hitinput_idspositionsslot_mappingc₈832c₉933c₁₀1034a₀00a₁11a₂22a₃33a₄424a₅525b₄416b₅517b₆61803912cu_seqlens_qkeys each query sees (seqlen_k)seq C: k = 11seq A: k = 6seq B: k = 7cu_seqlens_k = [0, 11, 17, 24]LM head keeps rows cu_seqlens_q[1:] − 1 = [2, 8, 11] (ringed)cu_seqlens_k[-1] = 24 > cu_seqlens_q[-1] = 12 → block_tables passed → K/V read from the paged cachemax_seqlen_q = 6max_seqlen_k = 11block_size = 4 (toy)bt C = [2, 5, 8]bt A = [0, 6]bt B = [3, 4]
    A prefill batch is one flat run of tokens, no padding. 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 at end − 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. Otherwise None, and attention uses the freshly computed k and v.

    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.

    nanovllm/engine/model_runner.pyL129–170GitHub ↗
    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, positions

    7.6prepare_decode: one token per sequence

    • input_ids is each sequence's last_token, the token sampled in the previous step.
    • positions is len(seq) − 1, the position that token occupies.
    • context_lens is len(seq). It includes the new token, whose K/V the store kernel writes in this same forward before attention reads the cache.
    • slot_mapping is the last block's base slot plus last_block_num_tokens − 1.
    • Block tables are always passed; decode always reads K/V from the cache.
    nanovllm/engine/model_runner.pyL172–188GitHub ↗
    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, positions

    7.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.

    nanovllm/engine/model_runner.pyL123–127GitHub ↗
    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_tables
    nanovllm/engine/model_runner.pyL190–193GitHub ↗
    190    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 temperatures

    7.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_eager is 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 first bs rows of the static buffers, and replay. slot_mapping is filled with −1 first, so padded rows store nothing (the Triton kernel returns early on −1), and context_lens is zeroed, so padded rows attend to nothing. block_tables is not cleared; stale entries are harmless because context_lens bounds what is read.
    • The LM head runs after the replay on only the real bs rows, so padding never pays for the 151,936-wide vocabulary projection.
    nanovllm/engine/model_runner.pyL195–220GitHub ↗
    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_ids

    7.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 (width ceil(max_model_len / block_size), 16 by default) and outputs ([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 Attention and 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.
    nanovllm/engine/model_runner.pyL222–257GitHub ↗
    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        )

    Which graph does a decode batch replay?

    The bucket list and the lookup in run_model.

    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.

    nanovllm/engine/model_runner.pyL50–59GitHub ↗
    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.

    nanovllm/utils/context.pyL1–27GitHub ↗
    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()
    FieldPrefillDecodeRead by
    is_prefillTrueFalseAttention, ParallelLMHead
    cu_seqlens_q[B+1] int32Noneattention; LM head picks last rows
    cu_seqlens_k[B+1] int32Noneattention
    max_seqlen_q / _kints0attention
    slot_mapping[N] int32[B] int32Triton store kernel
    context_lensNone[B] int32flash_attn_with_kvcache
    block_tablesonly with cached contextalwaysboth 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

    input_ids [N] · positions [N]VocabParallelEmbeddingQwen3DecoderLayer × 28(hidden, residual) → (hidden, residual)RMSNorm(hidden, residual)hidden [N, 1024]ParallelLMHeadSampler (rank 0 only)token_ids [B] → .tolist()self.model(input_ids, positions)captured as a CUDA graphfor decode, bs ≤ 512self.model.compute_logits()self.sampler(logits, temps)[N] → [N, 1024] (vocab 151,936 rows, ÷ tp)every layer gets positions for RoPEqkv [N, 4096] · MLP inner [N, 3072]adds the last residual, then normalisesprefill: keep the last row of each sequence[B, 1024] → logits [B, 151936]logits / T → softmax → argmax(p / Exp(1))
    The whole model, with Qwen3-0.6B sizes (hidden 1024, 28 layers, vocab 151,936). N is the number of tokens in the step: the packed prompt tokens in prefill, one per sequence in decode. B is the number of sequences. The framed part is 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.jsonValueWhere it matters
    hidden_size1024width of the residual stream
    num_hidden_layers28decoder layers, KV cache layers
    num_attention_heads / num_key_value_heads16 / 8GQA: two query heads per KV head
    head_dim128larger than 1024 / 16 = 64, so it must be read from the config
    intermediate_size3072MLP width
    vocab_size151,936embedding rows, logits width
    max_position_embeddings40,960RoPE table length; upper clamp for max_model_len
    rope_theta1,000,000RoPE base
    attention_biasfalseselects the QK-norm path
    tie_word_embeddingstrueLM head shares the embedding matrix
    rms_norm_eps1e-6every RMSNorm
    torch_dtypebfloat16default dtype while building the model

    9.2Qwen3ForCausalLM: the outer shell

    • packed_modules_mapping tells the loader that checkpoint tensors named q_proj, k_proj, v_proj go into slices of the fused qkv_proj, and gate_proj, up_proj into slices of gate_up_proj (Step 11).
    • With tied embeddings, lm_head.weight.data is pointed at the embedding's storage. One matrix serves both, so it is stored once.
    • forward stops at the final norm and returns hidden states; compute_logits is a separate call. The runner relies on that split: only forward is captured in CUDA graphs.
    nanovllm/models/qwen3.pyL186–216GitHub ↗
    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.

    nanovllm/models/qwen3.pyL162–183GitHub ↗
    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_states

    9.4Qwen3DecoderLayer: fused residual adds

    Qwen3DecoderLayer iresidualhidden =MLP out of i−1(not yet added)input_layernorm(h, r)r ← h + rh ← norm(r)post_attention_layernormr ← h + rh ← norm(r)self_attnmlpr (bf16)r outh out
    Pre-norm transformer with the residual add deferred. No layer ever computes 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 getattr defaults so the same class can load Qwen2-style configs: attention_bias defaults to True (Qwen2 has QKV bias), head_dim to None (then derived), rope_theta to 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_layernorm does that.
    nanovllm/models/qwen3.pyL120–159GitHub ↗
    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, residual

    9.5Qwen3Attention: projections, QK-norm, RoPE

    hidden [N, 1024]qkv_proj[N, 4096]q [N, 16, 128]k [N, 8, 128]v [N, 8, 128]q_normk_normRoPEpositionsstore_kvcacheTriton · slot_mappingk_cache / v_cacheflash-attnvarlen (prefill)with_kvcache (decode)paged K/V via block_tableo [N, 16, 128] → flatten(1, −1) → [N, 2048] → o_proj (row-parallel) → [N, 1024]
    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_size and kv_size are per-rank widths used to split the fused projection output.
    • scaling = head_dim ** -0.5 is the softmax scale passed to FlashAttention.
    • o_proj maps total_num_heads × head_dim (2048) back to hidden_size (1024) and is row-parallel, so it ends with the block's single all_reduce.
    • If rope_scaling is a dict, only its rope_theta is read. No RoPE scaling method (YaRN, linear, dynamic) is implemented.
    • get_rope is memoised, so all 28 layers share one RotaryEmbedding and one cos/sin table.
    • q_norm and k_norm exist 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].
    nanovllm/models/qwen3.pyL14–88GitHub ↗
    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 output

    9.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.

    nanovllm/models/qwen3.pyL91–117GitHub ↗
    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 x
    nanovllm/models/qwen3.pyL1–11GitHub ↗
    1import 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, ParallelLMHead

    Step 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

    column-parallel: split output rows, no communicationrow-parallel: split input columns, sum the partialsrank 0x [N, 1024]gate_up_projrows: gate[0:1536]+ up[0:1536] → [N, 3072]SiluAndMul→ [N, 1536]down_projcols [0:1536]→ partial [N, 1024]rank 1x [N, 1024]gate_up_projrows: gate[1536:3072]+ up[1536:3072] → [N, 3072]SiluAndMul→ [N, 1536]down_projcols [1536:3072]→ partial [N, 1024]all_reducesum over ranksy on every rankqkv_proj.weight rows on one rank (tp = 2): local heads only, in q | k | v orderq · 8 heads × 128 = 1024 rowsk · 4 × 128 = 512v · 4 × 128 = 512
    Megatron-style tensor parallelism on the Qwen3-0.6B MLP with two GPUs. Because 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).

    nanovllm/layers/linear.pyL1–34GitHub ↗
    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
    ClassShardsCommunicationUsed for
    ReplicatedLinearnothingnonenot used by Qwen3; available for other models
    ColumnParallelLinearoutput rows (tp_dim = 0)none; output stays shardedbase class of the next two
    MergedColumnParallelLineareach of several stacked outputs separatelynonegate_up_proj
    QKVParallelLinearquery heads and KV heads separatelynoneqkv_proj
    RowParallelLinearinput columns (tp_dim = 1)all_reduce of the outputo_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.

    nanovllm/layers/linear.pyL37–73GitHub ↗
    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.

    nanovllm/layers/linear.pyL76–93GitHub ↗
    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.

    nanovllm/layers/linear.pyL96–128GitHub ↗
    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.

    nanovllm/layers/linear.pyL131–156GitHub ↗
    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 y

    10.2embed_head.py: vocab-parallel embedding and LM head

    token idsmasklocal indexF.embedding · maskall_reduce[17, 98000, 5]rank 0 owns [0, 75968)[1, 0, 1][17, 0, 5][e₁₇, 0, e₅][17, 98000, 5]rank 1 owns [75968, 151936)[0, 1, 0][0, 22032, 0][0, e₉₈₀₀₀, 0]sum of partials[e₁₇, e₉₈₀₀₀, e₅]on every rankParallelLMHead reuses the same sharding: each rank computes logits for its own 75,968 ids,dist.gather sends them to rank 0, which concatenates them into [B, 151936]; the others get None.
    Embedding with tp = 2 on Qwen3's 151,936-token vocabulary. Out-of-range ids are mapped to local row 0 so the lookup is always legal, and their rows are zeroed afterwards. Exactly one rank contributes a non-zero row per token, so the sum is the true embedding.

    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.

    nanovllm/layers/embed_head.pyL1–42GitHub ↗
    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 y

    ParallelLMHead 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.

    nanovllm/layers/embed_head.pyL45–66GitHub ↗
    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 logits

    10.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.

    nanovllm/layers/attention.pyL1–30GitHub ↗
    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.

    nanovllm/layers/attention.pyL33–40GitHub ↗
    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_cache and v_cache start as empty tensors; allocate_kv_cache replaces 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_func on the fresh q, k, v with cu_seqlens describing the packed sequences.
    • Prefill with cached context (block_tables set): the same function, but k and v are the whole paged caches and block_table tells it where each sequence's keys are. With causal=True and 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_kvcache with q reshaped to [B, 1, heads, head_dim], reading context_lens keys per sequence through the block table.
    • GQA needs no special handling; FlashAttention accepts fewer KV heads than query heads.
    nanovllm/layers/attention.pyL43–75GitHub ↗
    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 o

    10.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_forward adds the residual in float32, keeps the sum, cast to model dtype, as the new residual, and normalises it. Under @torch.compile the 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.
    nanovllm/layers/layernorm.pyL1–50GitHub ↗
    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 cos and sin for every position up to max_position_embeddings and 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=False keeps it out of the state dict.
    • forward gathers rows by positions and rotates q and k. Positions come from the runner, so chunked prefill and prefix hits get the right angles with no extra work.
    • apply_rotary_emb pairs element j with element j + d/2 (the rotate-half layout used by GPT-NeoX and Hugging Face Qwen), not adjacent elements. It computes in float32.
    • rotary_dim must equal head_size; partial rotary is not supported.
    • @lru_cache(1) on get_rope returns the same module for identical arguments, so the table is built once per model.
    nanovllm/layers/rotary_embedding.pyL1–59GitHub ↗
    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_emb

    10.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.

    nanovllm/layers/activation.pyL1–11GitHub ↗
    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) * y

    10.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.

    nanovllm/layers/sampler.pyL1–12GitHub ↗
    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_tokens

    Temperature and the exponential race

    Six made-up logits for the token after “The capital of France is”.

    Bars show softmax(logits / T).

    Step 11Weight loading: utils/loader.py

    Twenty-eight lines map Hugging Face checkpoint names onto nano-vllm's fused, sharded parameters.

    1. Read the model's packed_modules_mapping (empty for models without one).
    2. Open every *.safetensors file in the model directory on the CPU and iterate over its tensor names.
    3. If a mapping key is a substring of the name, rewrite the name to the fused parameter, fetch it, and call its weight_loader with the shard id. break stops at the first matching key.
    4. The for … else branch runs when no key matched: fetch the parameter by its own name and use its weight_loader, or default_weight_loader (a plain copy) when it has none, as RMSNorm weights do.
    nanovllm/utils/loader.pyL1–28GitHub ↗
    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 tensorParameterLoader
    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 / 1MergedColumnParallelLinear.weight_loader
    …self_attn.o_proj.weight, …mlp.down_proj.weightsame nameRowParallelLinear.weight_loader
    model.embed_tokens.weight, lm_head.weightsame nameVocabParallelEmbedding.weight_loader
    …q_norm.weight, …input_layernorm.weight, model.norm.weightsame namedefault_weight_loader
    • model.get_parameter raises if a checkpoint tensor has no matching parameter, so unexpected tensors fail loudly. Parameters missing from the checkpoint are not detected; they keep whatever torch.empty left in memory.
    • Qwen3-0.6B ships one model.safetensors of 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-value lm_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 with copy_.

    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

    StageCodeWhat happens
    ScheduleScheduler.scheduleNo cache hits. A needs 2 blocks, B needs 1. 350 tokens fit the 16,384 budget, so both become RUNNING.
    Inputsprepare_prefillinput_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.
    Pathrun_modelPrefill is always eager.
    EmbedVocabParallelEmbedding[350] → [350, 1024]
    Per layer × 28Qwen3DecoderLayerqkv_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 normQwen3Model.norm[350, 1024]
    LM headParallelLMHeadRows cu_seqlens_q[1:] − 1 = [299, 349] → [2, 1024] → logits [2, 151936].
    SampleSamplertemperatures [2] → token ids [2] → .tolist()
    BookkeepingpostprocessA'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

    StageCodeWhat happens
    ScheduleScheduler.scheduleNothing waiting. None of 300, 1000, 57 is ≡ 1 mod 256, so no new blocks. Three sequences, one token each.
    Inputsprepare_decodeinput_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).
    Pathrun_modelbs = 3 → graph for 4. Rows 0–2 get real data; row 3 gets slot −1 and context length 0.
    Inside the graphQwen3ForCausalLM.forwardSame 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 headcompute_logits(outputs[:3])No row selection in decode: [3, 1024] → [3, 151936].
    SampleSamplerThree token ids.
    BookkeepingpostprocessNo 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.

    TechniqueWhereMechanismWhat it saves
    Paged KV cacheBlockManager, AttentionFixed 256-token blocks addressed through block tablesFragmentation; lets sequences grow without moving data
    Prefix cachingcompute_hash, can_allocate, hash_blocksChained xxh64 per full block; shared blocks are ref-countedRecomputing shared prompt prefixes and memory to hold them twice
    Lazy evictiondeallocate, _allocate_blockFreed blocks keep their hash until recycled, least recently freed firstTurns finished requests into a cache for future ones
    Chunked prefillscheduleThe first sequence of a batch may be cut to the token budgetLets prompts longer than max_num_batched_tokens run at all
    Continuous batchingSchedulerThe batch is rebuilt every stepIdle GPU time waiting for the slowest request
    Recompute preemptionpreemptFree a victim's blocks and re-prefill it later, mostly from cacheDeadlock when memory runs out, with no CPU swap space
    Varlen packingprepare_prefillcu_seqlens instead of paddingFLOPs on padding tokens
    Last-token logitsParallelLMHeadSelect one row per sequence before the vocab matmulA [tokens × 151,936] matmul in prefill
    CUDA graphscapture_cudagraph, run_model36 decode graphs, padded to the nearest bucketCPU launch overhead of hundreds of small kernels per decode step
    torch.compileRMSNorm, RoPE, SiluAndMul, SamplerElementwise chains fused into generated Triton kernelsMemory round-trips between elementwise ops
    Fused residual + normQwen3DecoderLayeradd_rms_forward does the add and the norm togetherA separate read and write of the residual stream per sublayer
    Fused projectionsQKVParallelLinear, MergedColumnParallelLinearq, k, v in one matmul; gate and up in one matmulKernel launches and repeated reads of the input
    Triton KV storestore_kvcache_kernelOne program per token writes K and V at its slotPython-level scatter into the paged cache
    Tensor parallelismlinear.py, embed_head.pyMegatron-style column/row split; two all-reduces per layerPer-GPU weight memory and per-GPU matmul time
    Shared-memory IPCwrite_shm / read_shm, __getstate__One pickle into a shared segment plus Events; decode sends only last tokensQueue/pipe copies and serialising whole token lists
    Pinned, async copiesprepare_*pin_memory=True then .cuda(non_blocking=True)Blocking host-to-device transfers
    Measured KV sizingwarmup_model, allocate_kv_cacheProfile the worst prefill, give the rest to the cacheGuesswork that wastes memory or runs out of it
    Exponential-race samplingSamplerargmax(p / Exp(1))torch.multinomial
    Tied embeddingsQwen3ForCausalLMLM head points at the embedding's storageA 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.

    nanovllm/engine/model_runner.pyL210GitHub ↗
    210            graph_vars["block_tables"][:bs, :context.block_tables.size(1)] = context.block_tables
    nanovllm/engine/scheduler.pyL71GitHub ↗
    71        assert scheduled_seqs

    14.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.

    nanovllm/engine/model_runner.pyL202GitHub ↗
    202            graph = self.graphs[next(x for x in self.graph_bs if x >= bs)]
    nanovllm/engine/model_runner.pyL234GitHub ↗
    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.

    nanovllm/engine/llm_engine.pyL21GitHub ↗
    21        Sequence.block_size = config.kvcache_block_size
    nanovllm/engine/model_runner.pyL181GitHub ↗
    181            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.
    • running can exceed max_num_seqs; only its first max_num_seqs sequences 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. generate returns dicts despite its list[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_dim must 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 exposes dtype on 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

    1. Watch the prefix cache. In Scheduler.schedule, log num_cached_blocks for each admitted sequence, then run example.py with two prompts that share a long system prompt.
    2. Add greedy decoding. Relax the assertion in SamplingParams, and in Sampler.forward use argmax on rows whose temperature is 0 before dividing.
    3. Add top-k. Carry a top_k field from SamplingParams through Sequence and prepare_sample, and mask logits below each row's k-th largest value before softmax.
    4. Validate request length. Reject or truncate in add_request when len(prompt) + max_tokens > max_model_len (14.1).
    5. Fix the bucket lookup. Append max_bs to graph_bs when it is not already the last bucket (14.2).
    6. Fix worker block size. Set Sequence.block_size = config.kvcache_block_size in ModelRunner.__init__ (14.3).