7.5 TensorRT-LLM, TGI & Other Serving Stacks¶

The serving stack you choose determines whether your model runs at 10 tokens/second on a laptop or 10,000 tokens/second on a data-center cluster. The transformer forward pass is the same mathematics in both cases — what differs is how aggressively the serving runtime compiles kernels, manages memory, schedules requests, and exposes hardware capabilities. This chapter maps the complete landscape: NVIDIA’s TensorRT-LLM for peak GPU throughput, HuggingFace’s Text Generation Inference (TGI) for portable production deployments, and llama.cpp / Ollama / LMDeploy / MLC for the long tail of hardware and use cases.

Before diving in, make sure you are comfortable with how the prefill and decode phases work and why the KV cache is central to serving economics — see The Anatomy of LLM Inference: Prefill, Decode & The KV Cache. Continuous batching and in-flight request scheduling, covered in Continuous Batching & Request Scheduling, underpin every production system we discuss here.

The Performance Hierarchy¶

Before benchmarking any serving framework, it helps to understand the ceiling imposed by hardware. The roofline model (see The Roofline Model & Performance Engineering) says decode throughput is bounded by memory bandwidth: each decode step reads the full set of model weights for every generated token. Take a 70 B model stored in FP16 (140 GB) sharded across two A100-80 GBs with ~2 TB/s of HBM bandwidth each. Every decode step streams all 140 GB — 70 GB out of each GPU’s HBM, in parallel — so one token takes at least \(70\ \text{GB} / 2\ \text{TB/s} = 35\ \text{ms}\), a ceiling near 29 tokens/second at batch size 1, before tensor-parallel all-reduces eat into it and before batching amortizes those same weight reads across concurrent requests.

Every framework we discuss is fighting to get close to this ceiling. The closer a system gets, the more it relies on compiled, fused, hardware-specific kernels — and the less portable it becomes. That tension is the through-line of this chapter.

TensorRT-LLM: The Compiled Engine Approach¶

What TensorRT-LLM Is¶

TensorRT-LLM (TRTLLM) is NVIDIA’s open-source library for building highly optimized TensorRT engines from transformer checkpoints. A TensorRT engine is a serialized, hardware-specific computation graph with all kernels pre-selected and fused at build time. Think of it as AOT (ahead-of-time) compilation for the GPU.

The key insight: PyTorch runs an interpreter that dispatches individual CUDA kernels at every step. TensorRT traces the full forward pass, applies a library of graph optimizations — layer fusion, constant folding, precision calibration — and emits a binary engine that the CUDA runtime can execute with minimal host overhead. For a model that runs millions of requests per day, eliminating Python-level dispatch overhead is material.

As of TensorRT-LLM 1.0 (2025), NVIDIA made a native PyTorch-based runtime the default flow and stabilized a high-level Python LLM API, so you no longer have to pre-compile an engine to get most of the performance. The AOT-engine path described below is still supported and still wins the last few percent of throughput, but much day-to-day work now runs through the PyTorch backend — which closes part of the deployment-friction gap with vLLM and TGI while keeping TensorRT-LLM’s tuned CUDA kernels.

Build and Run Pipeline¶

# Step 1: Convert a HuggingFace checkpoint to TensorRT-LLM format
# (here: Llama-2-7B in bfloat16, single GPU)
python tensorrt_llm/examples/llama/convert_checkpoint.py \
    --model_dir ./llama-2-7b-hf \
    --output_dir ./llama-2-7b-trtllm \
    --dtype bfloat16

# Step 2: Build the TensorRT engine
# max_batch_size and max_input_len are compile-time constants —
# choose them to match your production workload envelope.
# (Flag names drift between releases: recent trtllm-build versions fold
#  --max_input_len/--max_output_len into a single --max_seq_len. Always
#  check `trtllm-build --help` for the version you installed.)
trtllm-build \
    --checkpoint_dir ./llama-2-7b-trtllm \
    --output_dir ./llama-2-7b-engine \
    --max_batch_size 32 \
    --max_input_len 2048 \
    --max_output_len 512 \
    --use_inflight_batching \
    --paged_kv_cache enable \
    --gemm_plugin bfloat16 \
    --gpt_attention_plugin bfloat16

# Step 3: Run inference via the Python API
python -c "
import tensorrt_llm
from tensorrt_llm.runtime import ModelRunner
import tensorrt as trt

runner = ModelRunner.from_dir('./llama-2-7b-engine')
# Tokenize and run — runner handles batching and KV cache internally
outputs = runner.generate(
    batch_input_ids=[[1, 2, 3, 4, 5]],
    max_new_tokens=50,
)
print(outputs)
"

The --gemm_plugin and --gpt_attention_plugin flags activate NVIDIA’s hand-tuned GEMM and attention kernels, which are substantially faster than the TensorRT auto-scheduler can find on its own.

The 1.0 Path: the LLM API and trtllm-serve¶

On the PyTorch backend you skip Steps 1–2 entirely. The LLM class takes a HuggingFace checkpoint directly, and trtllm-serve puts an OpenAI-compatible endpoint in front of it — the same ergonomics as vllm serve, which is the point:

from tensorrt_llm import LLM, SamplingParams

# No convert_checkpoint.py, no trtllm-build: weights are loaded at startup
# and executed with TensorRT-LLM's tuned attention/GEMM kernels.
llm = LLM(model="meta-llama/Llama-3.1-8B-Instruct", tensor_parallel_size=1)
outputs = llm.generate(
    ["Explain the KV cache in one sentence."],
    SamplingParams(max_tokens=64, temperature=0.0),
)
print(outputs[0].outputs[0].text)

# Same thing as a server; speaks /v1/chat/completions and /v1/completions
trtllm-serve meta-llama/Llama-3.1-8B-Instruct --host 0.0.0.0 --port 8000

Reach for the AOT engine path when you have a fixed workload envelope and want the last few percent; reach for the LLM API when you are iterating on models or want a Triton-free deployment.

In-Flight Batching in TensorRT-LLM¶

TensorRT-LLM implements in-flight batching (also called continuous batching) through its Executor API. New requests are inserted into the active batch at the token boundary — pausing requests that have finished, inserting new ones, and continuing generation — without ever stopping the GPU. This is the same principle as vLLM (see vLLM: Architecture, PagedAttention & Internals), but TensorRT-LLM’s version is implemented inside a compiled TensorRT engine rather than in PyTorch.

# Minimal TensorRT-LLM Executor example (C++ API wrapped in Python)
# This illustrates the async request/response model.

from tensorrt_llm.executor import GenerationExecutor, GenerationRequest
import asyncio

async def serve_requests():
    """
    The Executor runs a background thread that continuously feeds the engine.
    Requests are submitted as GenerationRequest objects and picked up
    at the next scheduling interval (default: every decode step).
    """
    executor = GenerationExecutor.create(
        engine_dir="./llama-2-7b-engine",
        executor_config={
            "max_beam_width": 1,
            "scheduler_policy": "guaranteed_no_evict",  # vs "max_utilization"
        }
    )

    # Submit two requests concurrently — they will be batched automatically
    req_a = GenerationRequest(
        input_token_ids=[1, 234, 567],
        max_new_tokens=100,
        streaming=True,
    )
    req_b = GenerationRequest(
        input_token_ids=[1, 890, 123, 456],
        max_new_tokens=50,
        streaming=True,
    )

    executor.submit(req_a)
    executor.submit(req_b)

    # Stream tokens as they arrive
    async for token in req_a.aiter_tokens():
        print(f"A: {token}", end=" ", flush=True)

asyncio.run(serve_requests())

Paged KV Cache and Memory Management¶

TensorRT-LLM implements a paged KV cache that operates analogously to vLLM’s PagedAttention: KV tensors are stored in fixed-size “blocks” (on the order of 16 or 32 tokens per block), and a block manager allocates/frees blocks as sequences grow and terminate. This means a server can hold many more concurrent sequences than naïve pre-allocation would allow.

The key parameter is --kv_cache_free_gpu_mem_fraction (default 0.9): the fraction of free GPU memory TensorRT-LLM may use for KV blocks after the model weights are loaded.

Quantization in TensorRT-LLM¶

TensorRT-LLM has first-class support for INT8 weight-only quantization, INT8 SmoothQuant, FP8 (on H100/H200), NVFP4 (a 4-bit floating-point format with native tensor-core support on Blackwell B200/GB200, calibrated via NVIDIA ModelOpt), and GPTQ/AWQ. These are configured at engine build time:

# FP8 engine for H100 — uses calibration data to determine per-tensor scales
trtllm-build \
    --checkpoint_dir ./llama-2-7b-trtllm \
    --output_dir ./llama-2-7b-fp8-engine \
    --strongly_typed \
    --use_fp8_context_fmha enable \
    --max_batch_size 64 \
    --max_input_len 4096 \
    --max_output_len 1024

Because quantization scales are baked into the engine at build time, there is zero overhead at serving time — the quantized GEMM kernels simply execute with pre-computed scales.

For a deeper treatment of the quantization formats, see Quantization II: INT4/INT8/FP8, GGUF, bitsandbytes & QAT.

Multi-GPU Tensor Parallelism¶

TensorRT-LLM supports tensor parallelism and pipeline parallelism at build time. Specify --tp_size 4 (for 4-way tensor parallel) at the convert_checkpoint.py and trtllm-build stages; the library handles the all-reduce communication using NCCL. At inference time, the executor launches one process per GPU and co-ordinates automatically. See Multi-GPU & Multi-Node Inference for the broader parallelism strategies.

The Triton Inference Server Integration¶

NVIDIA’s long-standing production path is to serve TensorRT-LLM engines via Triton Inference Server using the tensorrtllm_backend. Triton handles HTTP/gRPC front-end, dynamic batching at the server layer, health checks, and metrics. TensorRT-LLM handles the GPU-side scheduling and execution.

Two things have changed by 2026. For a single-node deployment, trtllm-serve (above) removes the need for Triton at all. For multi-node fleets, NVIDIA Dynamo has become the orchestration layer of choice: it sits above the engine — TensorRT-LLM, vLLM, or SGLang — and adds KV-aware request routing, disaggregated prefill/decode pools, and tiered KV offload to CPU/SSD. See Disaggregated Prefill/Decode & Chunked Prefill.

TensorRT-LLM: What You Give Up¶

Everything in this subsection is the price of the AOT engine path; the PyTorch-backend LLM/trtllm-serve flow above pays none of it (and gives up a few percent of throughput in exchange). The engine is platform-specific and must be rebuilt for each GPU SKU (A100 vs H100 vs L40S). Build times for large models can be 30–60 minutes. The max_batch_size and max_input_len are fixed at build time — you cannot exceed them at runtime. If your traffic suddenly shifts to very long prompts, you need a different engine. These constraints mean TensorRT-LLM is most suitable for dedicated GPU fleets with predictable workloads.

Text Generation Inference (TGI)¶

Architecture Overview¶

Hugging Face’s Text Generation Inference is a Rust-based server with a Python model-runner process. The Rust front-end handles HTTP/gRPC routing and request queuing; the Python process runs the model using PyTorch + custom CUDA kernels. Importantly, TGI ships its own custom attention kernels (including a FlashAttention-based implementation) and its own continuous batching scheduler.

As of 2026, Hugging Face has moved TGI into maintenance mode — the repository was archived in March 2026 — and now steers new production work toward vLLM, SGLang, and llama.cpp/MLX, the very engines that adopted TGI’s transformers-native architecture. TGI still runs well, and its multi-backend support can even front a vLLM or TensorRT-LLM engine behind the same API, but for a greenfield 2026 deployment vLLM is the more common default. The architecture below remains an instructive template for what a production LLM server actually does.

Running TGI¶

# Launch TGI with Docker — NVIDIA GPU required for full throughput
docker run --gpus all \
    -p 8080:80 \
    -v $PWD/model_cache:/data \
    ghcr.io/huggingface/text-generation-inference:2.4 \
    --model-id meta-llama/Llama-3-8B-Instruct \
    --quantize bitsandbytes-nf4 \
    --max-input-tokens 4096 \
    --max-total-tokens 6144 \
    --max-batch-prefill-tokens 16384 \
    --num-shard 1  # tensor parallel degree

# Send a request — TGI implements the Messages API compatible with OpenAI
curl http://localhost:8080/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "tgi",
    "messages": [{"role": "user", "content": "Explain KV cache in one paragraph."}],
    "max_tokens": 200,
    "stream": true
  }'

TGI exposes a /generate endpoint (its native API), a /v1/chat/completions OpenAI-compatible endpoint, and a /v1/completions endpoint. The OpenAI-compatible surface makes it easy to drop TGI behind any client that speaks the OpenAI protocol.

TGI’s Continuous Batching Scheduler¶

TGI implements a “token budget” scheduler: each request is assigned a max_total_tokens budget, and requests are grouped into batches where the total token budget fits within the KV cache. The scheduler preempts requests if memory pressure grows (evicting KV blocks and later recomputing them). The Rust router enforces these constraints before passing requests to the Python runner, which means request validation and queueing happen with very low latency.

# Simplified pseudocode illustrating TGI's waiting queue logic
# Real implementation is in Rust + a Python model server side

from dataclasses import dataclass

@dataclass
class Request:
    current_length: int
    max_total_tokens: int

class TGIScheduler:
    def __init__(self, max_batch_total_tokens: int):
        self.max_batch_total_tokens = max_batch_total_tokens
        self.waiting: list[Request] = []
        self.running: list[Request] = []

    def schedule(self) -> list[Request]:
        """
        Called every decode step. Fill the running batch up to the token budget.
        New requests are added from waiting if budget permits.
        Running requests keep their slot as long as they haven't finished.
        """
        budget_used = sum(r.current_length for r in self.running)
        for req in list(self.waiting):
            needed = req.max_total_tokens  # pre-allocated worst case
            if budget_used + needed <= self.max_batch_total_tokens:
                self.running.append(req)
                self.waiting.remove(req)
                budget_used += needed
        return self.running

TGI vs TensorRT-LLM: Key Differences¶

llama.cpp: The CPU/Edge Champion¶

What llama.cpp Does¶

Georgi Gerganov’s llama.cpp (2023) proved that a quantized transformer can run usefully fast on commodity hardware — a laptop CPU, an Apple M-series chip, or a consumer GPU. The project is a single C++ codebase with no external deep-learning dependencies that achieves high throughput through:

1. GGUF quantization: 2-bit through 8-bit per-channel integer quantization with mixed precision (Q4_K_M, Q5_K_M, etc.) that keeps key tensors at higher precision.
2. Highly optimized BLAS routines: AVX2/AVX-512 vectorized GEMM for x86, ARM NEON/SVE for ARM, Metal for Apple Silicon, and CUDA for NVIDIA GPUs — all in the same binary via compile-time backends.
3. mmap model loading: The model file is memory-mapped, so the OS controls paging. On a machine with enough RAM, the model loads in seconds; with limited RAM, the OS pages in only the layers currently needed.

GGUF Format¶

GGUF (GPT-Generated Unified Format) stores model weights, tokenizer data, and metadata in a single binary file. Quantized weights use a block quantization scheme: weights are grouped into blocks of 32 values, each block stores a float32 scale factor and the quantized integers. For Q4_K_M, 4-bit integers are stored for most weights with 6-bit integers for sensitive layers (attention projections, etc.):

where scale is a per-32-element float32. The storage cost for Q4_K_M is approximately \(4.5\) bits per weight after accounting for the scale overhead.

# Estimate GGUF model file size from parameter count
def gguf_size_gb(params_billions: float, bits_per_weight: float = 4.5) -> float:
    """
    Rough size estimate for a GGUF-quantized LLM.
    bits_per_weight: 4.5 for Q4_K_M, 5.5 for Q5_K_M, 8.5 for Q8_0
    """
    params = params_billions * 1e9
    bytes_per_weight = bits_per_weight / 8.0
    # Embeddings and norms are typically kept in FP16, ~5% of params
    emb_fraction = 0.05
    size_bytes = (
        params * (1 - emb_fraction) * bytes_per_weight
        + params * emb_fraction * 2.0  # FP16
    )
    return size_bytes / (1024**3)

# Llama-3-70B at Q4_K_M:
print(f"70B Q4_K_M: {gguf_size_gb(70, 4.5):.1f} GB")   # ~41 GB
print(f"70B Q5_K_M: {gguf_size_gb(70, 5.5):.1f} GB")   # ~49 GB
print(f"13B Q4_K_M: {gguf_size_gb(13, 4.5):.1f} GB")   # ~7.7 GB

Running llama.cpp Directly¶

The repository builds a set of small binaries; three of them are the whole workflow from a trained checkpoint to a served endpoint:

# Build (CUDA backend; omit -DGGML_CUDA for a CPU/Metal-only build)
cmake -B build -DGGML_CUDA=ON && cmake --build build -j

# Convert an HF checkpoint to GGUF, then quantize it
python convert_hf_to_gguf.py ./my-model-hf --outfile my-model-f16.gguf --outtype f16
./build/bin/llama-quantize my-model-f16.gguf my-model-q4_k_m.gguf Q4_K_M

# Serve an OpenAI-compatible endpoint with 4 concurrent slots, all layers on GPU
./build/bin/llama-server -m my-model-q4_k_m.gguf --port 8080 --parallel 4 -ngl 99

llama-server does continuous batching and prompt-prefix reuse across its slots, but its memory model is not paged: the context budget is partitioned statically across --parallel slots, so a slot’s usable context is roughly --ctx-size / --parallel and one long request cannot borrow blocks from an idle neighbour. That is the concrete price of not having a PagedAttention-style block manager, and it is why concurrency scales far worse here than in vLLM or TensorRT-LLM even when single-stream speed is competitive.

This same three-command path is exactly how Stack-100M gets from a trained checkpoint to a laptop demo — see Evaluation & Serving: Honest Benchmarks, int4 Quantization, and Running on a Laptop, which also covers the pre-tokenizer-hash trap that silently corrupts the GGUF of a model whose tokenizer the converter has never seen.

Ollama: A Developer-Friendly Frontend¶

Ollama wraps llama.cpp in a REST server with a model registry, automatic download, and a simple CLI. It is the fastest path from “zero” to running a local LLM. (Since 2025 Ollama has also shipped its own GGML-based engine for newer and multimodal architectures, so it no longer routes every model through llama.cpp — though llama.cpp still powers much of the local-inference ecosystem, including tools like LM Studio.)

# Install and run Llama-3-8B locally
ollama pull llama3.1:8b
ollama run llama3.1:8b "Explain transformer attention in two sentences."

# OpenAI-compatible API (on port 11434 by default)
curl http://localhost:11434/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "llama3.1:8b",
    "messages": [{"role": "user", "content": "What is speculative decoding?"}]
  }'

Ollama automatically detects available hardware (CUDA, Metal, CPU) and offloads as many layers as possible to GPU. The key flag is --num-gpu (or OLLAMA_NUM_GPU in the environment), which sets how many transformer layers run on GPU versus CPU.

Performance Profile¶

On Apple M3 Max with 128 GB of unified memory, a Llama-3-70B Q4_K_M model (≈41 GB) fits entirely in unified memory and typically achieves 20–30 tokens/second — genuinely useful for local development. On a consumer NVIDIA 4090 (24 GB), the 13B Q4_K_M fits fully in VRAM and achieves 80–100 tokens/second. For anything requiring production throughput (hundreds of tokens/second, many concurrent users), llama.cpp is not the right tool.

LMDeploy: Efficient Serving with TurboMind¶

LMDeploy, developed by Shanghai AI Laboratory (the InternLM team), provides a high-throughput serving stack that sits between TGI (easy to use) and TensorRT-LLM (maximum performance). Its core inference engine is TurboMind, a C++/CUDA implementation with:

• Custom blocked KV cache (similar in concept to PagedAttention)
• Continuous batching
• INT4 and INT8 quantization via AWQ
• A Python/gRPC serving layer

# Install and convert a model to TurboMind format
pip install lmdeploy

# Convert HF checkpoint to TurboMind internal format
lmdeploy convert llama3 /path/to/llama-3-8b-hf \
    --dst_path ./llama3-turbomind

# Launch the serving API
lmdeploy serve api_server ./llama3-turbomind \
    --server_port 23333 \
    --tp 1 \
    --cache_max_entry_count 0.8

# Or use the Python API directly
python -c "
from lmdeploy import pipeline, TurbomindEngineConfig

cfg = TurbomindEngineConfig(tp=1, cache_max_entry_count=0.8)
pipe = pipeline('./llama3-turbomind', backend_config=cfg)
response = pipe(['Hello! Explain TurboMind in one sentence.'])
print(response[0].text)
"

LMDeploy also supports the PyTorch backend (--backend pytorch) which is more portable but slower, mirroring TGI’s design philosophy. The TurboMind backend is specifically optimized for NVIDIA GPUs and resembles TensorRT-LLM in its use of custom CUDA kernels — but without the full AOT compilation step, making model loading much faster.

MLC-LLM: Universal Compilation via TVM¶

MLC-LLM (Machine Learning Compilation for LLMs), from the MLC team led by Tianqi Chen, applies Apache TVM’s compilation machinery to LLM serving. The goal is one codebase that runs on essentially any hardware: NVIDIA CUDA, AMD ROCm, Apple Metal, Intel Vulkan, ARM CPU, and even WebGPU (in the browser).

How MLC Differs¶

Where TensorRT-LLM compiles to a TensorRT engine (NVIDIA-specific), MLC compiles to TVM’s intermediate representation (TIR), then lowers TIR to hardware-specific code via TVM’s code generation backends. The compilation includes automatic schedule search (AutoTIR) that tunes GEMM tile sizes and memory layouts for the target device.

# Compile a model for Metal (Apple Silicon) using MLC
import mlc_llm
from mlc_llm import MLCEngine

# mlc_llm.build() emits a compiled library (.dylib on macOS, .so on Linux)
# This step is slow once; subsequent loads are fast.
mlc_llm.build(
    model="HF://meta-llama/Llama-3.2-3B-Instruct",
    target="apple/m3-gpu",   # or "cuda", "rocm", "vulkan", "webgpu"
    quantization="q4f16_1",  # 4-bit weights, FP16 activations
)

# Run inference — identical Python API regardless of hardware backend
engine = MLCEngine("./dist/Llama-3.2-3B-Instruct-q4f16_1-MLC")
response = engine.chat.completions.create(
    messages=[{"role": "user", "content": "What is PagedAttention?"}],
    max_tokens=100,
)
print(response.choices[0].message.content)

MLC is the right choice when you need to deploy to heterogeneous hardware (a fleet with a mix of NVIDIA, AMD, and Apple GPUs), or when you need to run in a browser via WebLLM (which compiles to WebGPU/WASM). The trade-off is that compilation is slow and the tuned performance on any single GPU is typically below TensorRT-LLM.

The Tradeoff Space: A Practitioner’s Framework¶

Decision Matrix¶

Portability vs Peak Performance: The Core Tension¶

The fundamental constraint is this: the more a runtime commits to a specific hardware target, the more aggressive its optimizations can be. TensorRT-LLM achieves peak performance by: 1. Selecting the exact GEMM algorithm for each matrix shape at build time (profile-guided kernel selection) 2. Fusing operations that PyTorch would run as separate kernels (e.g., attention + residual + layer norm in one pass) 3. Emitting machine code rather than dispatching at runtime

Every step away from that specificity costs performance. TGI’s PyTorch-based model runner dispatches CUDA kernels at runtime (some overhead) but works on any HuggingFace model on any CUDA GPU. llama.cpp’s AVX2 kernels work on CPUs but obviously cannot exploit tensor cores.

The quantitative gap is real but often overstated in marketing: well-tuned vLLM or TGI on A100s typically reaches 85–95% of TensorRT-LLM’s throughput at matched batch sizes, while being far simpler to deploy and maintain. The remaining 5–15% matters at scale (thousands of dollars per day) but is often dominated by other system costs (networking, load balancing, tokenization) at moderate scale.

Measuring It Yourself¶

Never take a number like “85–95%” — including this chapter’s — on faith for your workload. Every stack ships a load generator, and any OpenAI-compatible endpoint can be driven by a shared one:

# TensorRT-LLM: engine/backend throughput sweep over a prompt dataset
trtllm-bench --model meta-llama/Llama-3.1-8B-Instruct throughput --dataset prompts.jsonl

# vLLM: online serving benchmark against a running server (TTFT/ITL percentiles)
vllm bench serve --model meta-llama/Llama-3.1-8B-Instruct \
    --dataset-name sharegpt --request-rate 8

# NVIDIA GenAI-Perf: engine-agnostic; drives any OpenAI-compatible endpoint
# (trtllm-serve, vLLM, SGLang, TGI, llama-server) through the same harness
genai-perf profile -m meta-llama/Llama-3.1-8B-Instruct \
    --endpoint-type chat --streaming --url localhost:8000

# llama.cpp: single-process prefill (-p) and decode (-n) microbenchmark
llama-bench -m my-model-q4_k_m.gguf -p 512 -n 128

Three rules make such comparisons honest. Fix the input and output length distribution across stacks (a stack that emits shorter completions “wins” on tokens/second for free). Sweep request rate, not batch size — production systems are open-loop, and the interesting number is the arrival rate at which p99 TTFT crosses your SLO. And report TTFT and inter-token latency as percentiles, since the eviction and preemption policies discussed above show up almost entirely in the tail. The cost model that turns these curves into dollars per million tokens is Inference Economics: Latency, Throughput & Cost. (Benchmark CLIs move fast; check --help for the version you installed.)

Kernel Highlights: What Makes These Engines Fast¶

Fused Attention Kernels¶

All production serving stacks today incorporate FlashAttention or a variant. The key insight (Dao et al., 2022) is that the standard attention computation is memory-bandwidth bound, not compute bound, and by tiling the Q/K/V matrices to fit in on-chip SRAM, we avoid round-trips to HBM. See FlashAttention I: IO-Awareness & The Online Softmax for the full derivation.

For paged KV caches, standard FlashAttention (which assumes contiguous key/value tensors) must be modified. TensorRT-LLM, vLLM, and TGI all ship custom “paged FlashAttention” kernels that index into the block table at each attention step.

// Simplified pseudocode for paged attention kernel (CUDA C++)
// Each thread block handles one query head for one sequence in the batch.

__global__ void paged_attention_kernel(
    float* output,           // [num_seqs, num_heads, head_dim]
    const float* queries,    // [num_seqs, num_heads, head_dim]
    const float** kv_blocks, // array of pointers to KV blocks
    const int* block_table,  // [num_seqs, max_blocks_per_seq]
    int head_dim,
    int block_size,          // tokens per KV block (e.g., 16)
    int seq_len
) {
    int seq_id = blockIdx.x;
    int head_id = blockIdx.y;

    // Accumulator for online softmax
    float acc[HEAD_DIM] = {0.0f};
    float max_score = -1e9f;
    float sum_exp = 0.0f;

    // Iterate over KV blocks for this sequence
    for (int block_idx = 0; block_idx * block_size < seq_len; ++block_idx) {
        int physical_block = block_table[seq_id * MAX_BLOCKS + block_idx];
        // Load K from this block, compute QK^T, update online softmax...
        // (full implementation follows standard online softmax pattern)
    }

    // Write final attended value to output
    // ...
}

For a full treatment of how PagedAttention fits into the memory manager, see PagedAttention & KV-Cache Memory Management.

GEMM Kernel Tuning¶

Decode-phase matrix multiplications are “tall-and-skinny”: the weight matrix is large (e.g., \([4096 \times 16384]\) for an FFN layer), but the activation tensor is only \([\text{batch\_size} \times 4096]\) — often with batch size below 32 during decode. This shape is very different from training GEMM, where batch × sequence length gives a “fat” activation matrix.

TensorRT-LLM uses a profiling pass during engine build to time dozens of cuBLAS/cuBLASLt configurations against your specific shapes and selects the best one. This profile-guided selection is a major source of its throughput advantage over frameworks that use PyTorch’s general-purpose cuBLAS dispatch.

KV Cache Quantization¶

Beyond paging, modern runtimes also quantize the KV cache itself. INT8 KV cache is now standard in TensorRT-LLM and supported in TGI and LMDeploy. The quantization is applied per-head-per-token with a float scale factor:

This halves KV cache memory at the cost of a small accuracy reduction (typically under 0.5 pp on standard benchmarks for INT8).

Structured Output and Constrained Generation¶

All major serving stacks support constrained generation — forcing the model to produce valid JSON, follow a regex, or conform to a grammar. The mechanisms differ:

• TGI uses Outlines (a regex/grammar-to-finite-state-automaton library) to mask logits during sampling to only valid next tokens.
• vLLM integrates Outlines similarly.
• TensorRT-LLM requires a custom logit processor passed to the executor.
• llama.cpp has built-in GBNF (GGML BNF) grammar support.

For a full treatment of constrained generation mechanics, see Structured & Constrained Generation.

Stack Selection Cheatsheet¶

Further Reading¶

• TensorRT-LLM repository — NVIDIA, github.com/NVIDIA/TensorRT-LLM. The primary reference for engine build options, supported models, and the Executor API.
• Text Generation Inference repository — Hugging Face, github.com/huggingface/text-generation-inference. The Rust router source and Python model server are both readable and well-documented.
• Kwon et al., “Efficient Memory Management for Large Language Model Serving with PagedAttention,” SOSP 2023 — foundational paper for block-based KV cache management; the ideas are implemented in vLLM, TGI, TRT-LLM, and LMDeploy.
• Dao et al., “FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness,” NeurIPS 2022 — the attention kernel that every production stack now ships.
• llama.cpp repository — ggml-org (Georgi Gerganov), github.com/ggml-org/llama.cpp. The GGUF specification and all backend implementations are in this repo.
• MLC-LLM: Universal LLM Deployment — Chen et al., MLC team, github.com/mlc-ai/mlc-llm. Describes the TVM-based compilation pipeline and cross-platform targets.
• LMDeploy repository — Shanghai AI Laboratory, github.com/InternLM/lmdeploy. TurboMind engine design documentation and AWQ quantization integration.
• Aminabadi et al., “DeepSpeed-Inference: Enabling Efficient Inference of Transformer Models at Unprecedented Scale,” SC 2022 — covers multi-GPU inference kernels and the design philosophy behind high-throughput serving at scale.

Exercises¶

1. A team builds a TensorRT-LLM engine for Llama-2-7B on an A100 with --max_input_len 2048 --max_batch_size 32. Two things then happen: (a) a product change starts sending prompts of up to 6000 tokens, and (b) the fleet is upgraded from A100 to H100 GPUs. Explain, using this chapter’s description of what a TensorRT engine is, why each of these events forces a rebuild — and contrast this with how TGI would handle the same two events.

2. Using the chapter’s gguf_size_gb helper, estimate by hand the on-disk size of an 8B-parameter model quantized to Q4_K_M (4.5 bits/weight), assuming 5% of parameters (embeddings/norms) are kept in FP16. Show the arithmetic and give the answer in GB (with \(1\ \text{GB} = 1024^3\) bytes).

3. The chapter’s roofline argument says decode is memory-bandwidth bound: each decode step reads the full weight matrix once per token. For an A100 with 2 TB/s of HBM bandwidth, compute the batch-size-1 decode ceiling (tokens/sec) for an 8B model (a) in BF16 and (b) quantized to ~4.5 bits/weight. What does the ratio tell you about why weight quantization speeds up decode, and where does the ceiling stop improving as you raise the batch size?

4. You serve Llama-3-8B (32 layers, GQA with 8 KV heads, head dimension 128) on a single A100-80 GB in BF16. Following the chapter’s worked-example method, compute: (a) memory left for the KV cache after weights, (b) KV bytes per token across all layers, © the maximum number of concurrent 4096-token sequences you can hold, and (d) the new count if you enable INT8 KV-cache quantization. Use \(1\ \text{GB}=10^9\), \(1\ \text{KB}=10^3\) bytes, as the chapter does.

5. Turn the worked example into reusable code. Implement a function max_concurrent_seqs(hbm_gb, param_billions, num_layers, kv_heads, head_dim, ctx_len, weight_bytes=2, kv_bytes=2) that returns the number of concurrent sequences of length ctx_len that fit, in the chapter’s decimal-GB convention. Then call it to recompute the Exercise 4 scenario for both BF16 KV and INT8 KV, and reconcile the result with Exercise 4’s numbers.

def max_concurrent_seqs(
    hbm_gb: float,
    param_billions: float,
    num_layers: int,
    kv_heads: int,
    head_dim: int,
    ctx_len: int,
    weight_bytes: int = 2,   # BF16 weights
    kv_bytes: int = 2,       # 2 = BF16 KV, 1 = INT8 KV
) -> int:
    """Concurrent sequences that fit in HBM after loading weights.

    Decimal HBM sizing (1 GB = 1e9 bytes); KV cache in exact bytes.
    """
    hbm = hbm_gb * 1e9
    weights = param_billions * 1e9 * weight_bytes
    kv_free = hbm - weights
    if kv_free <= 0:
        return 0
    # 2x for K and V; per token, per layer, summed over layers.
    kv_per_token = 2 * kv_heads * head_dim * kv_bytes * num_layers
    max_tokens = kv_free / kv_per_token
    return int(max_tokens // ctx_len)

# Llama-3-8B on one A100-80GB, 4096-token sequences
bf16 = max_concurrent_seqs(80, 8, 32, 8, 128, 4096, kv_bytes=2)
int8 = max_concurrent_seqs(80, 8, 32, 8, 128, 4096, kv_bytes=1)
print(f"BF16 KV: {bf16} sequences")   # BF16 KV: 119 sequences
print(f"INT8 KV: {int8} sequences")   # INT8 KV: 238 sequences

6. TensorRT-LLM’s executor exposes two scheduler policies: guaranteed_no_evict and max_utilization. Explain what each does when the KV cache fills up, and argue which one the Interview Corner’s low-TTFT customer-support deployment should use — and why the other policy would hurt its latency SLO even though it can pack more requests.