TensorRT-LLM: Compiling Inference Engines Ahead of Time

vLLM Architecture and SGLang and RadixAttention followed two engines that run a model through PyTorch and decide each step what to run next. TensorRT-LLM, NVIDIA’s engine, started from a different premise: compile the model once, ahead of time, into an engine, a serialized plan for one GPU model in which operations are fused into fewer kernels and every kernel has been timed and chosen for the shapes it will serve. For most of the project’s life, serving a model meant building its engine first.

That is no longer how TensorRT-LLM works. Since version 1.0 it runs models through PyTorch by default, and the 1.3 pre-releases since July 2026 remove the engine backend. This post measures the two things an engine build did, fusing and tuning, on a laptop GPU (TensorRT-LLM itself needs an NVIDIA GPU on Linux), then follows the rest of the design and the move away from it.

Table of contents

What an engine build did

TensorRT-LLM is NVIDIA’s open-source (Apache 2.0) LLM inference library, first released as v0.5.0 in October 2023. It took over from FasterTransformer, an earlier NVIDIA library, and was built on TensorRT, NVIDIA’s compiler for neural network inference. There is no paper for it, as there is for vLLM and SGLang; the sources are the repository, its release notes and its documentation.

The original workflow had a build step. A model was described with TensorRT-LLM’s Python API as a TensorRT network, and trtllm-build compiled it. TensorRT’s documentation describes what the builder does with a network: it “picks the fastest available GPU kernel (a low-level GPU function) for each layer in your network”, producing a “serialized, hardware-specific binary”. Three things happened in that step:

• Fusion. Chains of small operations, such as a normalization’s reductions and multiplies, or an activation followed by a product, became single kernels, often from hand-written plugins.
• Tuning. For each layer, candidate kernels were timed on the target GPU for the build’s shape profile, the range of batch sizes and sequence lengths the engine would accept, and the fastest kept.
• Precision. A quantized format and its scales were fixed in the engine.

The engine was the result: fast, and tied to the GPU model and TensorRT version it was built with and to its shape profile. A new GPU, a longer context or a bigger batch meant a new build. The next two sections measure the first two steps.

Fusing operations, measured

A GPU runs a model as a sequence of kernels. Each one reads its inputs from memory and writes its output back, and each one costs a launch. A small operation, such as adding a constant, does almost no arithmetic, so its cost is the trip through memory and the launch. Fusing a chain of them into one kernel reads the input once, keeps the intermediate values in registers, and writes once.

Qwen2.5-1.5B-Instruct has 28 layers, each with two RMSNorms and a gated MLP. As written in Hugging Face’s code, an RMSNorm is eight PyTorch operations (cast to float, square, mean, add epsilon, reciprocal square root, multiply, cast back, multiply by the weight), and the MLP computes down(silu(gate(x)) * up(x)): two matmuls into 8,960 columns, a SiLU and a multiply. TensorRT-LLM’s own modules fuse both: its gated MLP has one gate_up_proj matmul followed by a fused swiglu, and its RMSNorm adds the residual and normalizes in one kernel. Two similar fusions (without the residual add), written as Metal kernels for the laptop’s GPU (Metal is Apple’s GPU language, so these scripts run only on Apple silicon), are swapped into the model:

def fused_norm(self, x):
    global metal; metal += 1
    threads, vec = (512, 4) if x.numel() // x.shape[-1] <= 16 else (128, 4)
    return kernels.rmsnorm(x.contiguous(), self.weight, self.variance_epsilon, threads, vec)

def fused_mlp(self, x):
    global metal; metal += 1
    return self.down_proj(kernels.swiglu(F.linear(x, self.gate_up)))

Full script: fuse.py

self.gate_up is the two projection matrices stacked, so one matmul produces both halves. The SiLU-and-multiply kernel reads the two halves and writes one value per thread, rounding where PyTorch rounds:

kernel void swiglu(device const half* gu, device half* y, constant uint& n,
                   uint i [[thread_position_in_grid]]) {
    uint row = i / n, j = i % n;
    float g = gu[row * 2 * n + j];
    y[i] = half(g / (1 + exp(-g))) * gu[row * 2 * n + n + j];  // rounded where PyTorch rounds
}

Full script: kernels.py

The script counts the operations one decoded token runs (PyTorch’s, allocations included and views not, plus the Metal kernels), times greedy decoding of 128 tokens (Hugging Face’s generate, including its 64-token prompt) and a 2,048-token prefill (including the logits for every position), alternating the two versions over seven rounds so that any drift in the GPU’s speed hits both, and compares outputs. On an Apple M4:

as written: 1425 operations per decoded token;  39.7 ms per decoded token; prefill of 2,048 tokens 2289.3 ms
fused     : 1055 operations per decoded token;  38.0 ms per decoded token; prefill of 2,048 tokens 2121.3 ms
next-token logits over 512 positions: largest difference 0.051 (logits span -20 to 29); same top token at 510 of 512, where the top two differ by under 0.1 at 48
greedy output: the first 23 of 128 tokens are the same
weights read per decoded token: 3.09 GB; at 38.0 ms per token, 81 GB/s

Fusion removed 370 operations from every decoded token, 26%. Prefill, which pushes 2,048 rows through every operation, got 7% faster, since each removed normalization step was a pass over a 2,048 by 1,536 array, and the fused activation skips writing and rereading a 2,048 by 8,960 intermediate. Decoding got only 4% faster, and the last line suggests why: each decoded token reads all 3.09 GB of weights, and at 38 ms per token that is 81 GB/s, close to what this laptop’s memory delivers (the copy kernel in Paged Attention measured 84 GB/s). The weight reads set the pace, and the CPU can queue the next kernels while the GPU reads, so the operations removed save little time here. (The script does not measure the CPU’s side separately.)

On a data-center GPU the balance flips. An H100 SXM reads memory at 3.35 TB/s, about 40 times faster, so the same weights take under a millisecond, and issuing on the order of a thousand kernels from Python can take longer than the GPU needs to run them. The GPU then waits on the CPU. That is why TensorRT-LLM fuses so aggressively, and why it also captures CUDA graphs: a recording of a whole decoding step’s kernel launches that the GPU replays with one launch from the CPU.

The outputs are not bit-identical. The fused RMSNorm adds up a row in a different order, the merged matmul may accumulate differently, and Metal’s math functions round differently, so a value can differ in its last bits; the largest logit difference was 0.051. That is enough to change the top token where the two best candidates are nearly tied (two of 512 positions here), and once greedy decoding picks a different token, the rest of the output follows a different path: here after 23 tokens. Any change in how a model’s arithmetic is grouped can do this, so a fused engine’s output should be compared with the original’s up to numerical differences, not token for token.

Auto-tuning, measured

A kernel can be written many ways: how many threads share one row, how many numbers each thread reads at once, how work is split between groups of threads. No version is fastest for every shape. With one row, a version that gives the row few threads leaves most of the GPU idle; with thousands of rows, a version that gives each row many threads spends its time coordinating them. TensorRT’s builder settled this by measurement: time each candidate at each shape, keep the fastest.

The fused RMSNorm kernel above has two such choices: threads per row (32 to 1,024) and numbers per read (2 or 4), twelve versions in all. Each is timed at four shapes: one row (decoding one sequence), 16 rows (decoding 16), 256 rows (a chunk of a prompt) and 4,096 rows (a long prompt). The GPU is warmed up first, and every shape and version is run once per round for nine rounds, keeping the median:

runs = {(r, tv): [] for r in xs for tv in VERSIONS}
for round in range(9):                        # every shape and version once per round;
    for r, tv in runs:                        # the median over rounds is kept
        runs[r, tv].append(timed(lambda: kernels.rmsnorm(xs[r], w, EPS, *tv, out=ys[r])))
table = {r: {tv: sorted(runs[r, tv])[4] for tv in VERSIONS} for r in xs}
best = {r: min(table[r], key=table[r].get) for r in table}

Full script: tune.py

microseconds per call; rows: 1, 16, 256, 4096 (fastest marked *)
  32 threads x 2     9.0      8.8     10.3    285.8 
  32 threads x 4     6.6      6.2      8.3    289.3 
  64 threads x 2     6.5      6.4      9.6    299.3 
  64 threads x 4     5.6      5.1      8.3    284.9 
 128 threads x 2     5.5      5.3      9.6    285.1 
 128 threads x 4     4.7      4.3      8.2*   290.0 
 256 threads x 2     4.6      4.8     10.3    300.4 
 256 threads x 4     4.5      4.9      9.3    278.8 
 512 threads x 2     4.3      4.6     12.7    294.5 
 512 threads x 4     4.3      4.2*    11.0    275.1*
1024 threads x 2     4.5      5.5     16.6    288.9 
1024 threads x 4     4.2*     4.8     14.3    277.0 
   1 rows: fastest   4.2 us, slowest 2.1x as long; within 5% of the fastest: 512x2, 512x4, 1024x4
  16 rows: fastest   4.2 us, slowest 2.1x as long; within 5% of the fastest: 128x4, 512x4
 256 rows: fastest   8.2 us, slowest 2.0x as long; within 5% of the fastest: 32x4, 64x4, 128x4
4096 rows: fastest 275.1 us, slowest 1.1x as long; within 5% of the fastest: 32x2, 64x4, 128x2, 256x4, 512x4, 1024x4

The shapes want opposite things. One row wants 512 threads or more; 16 rows are within 5% of the best only at 128 and 512 threads reading four numbers at a time. With 32 threads reading two at a time, both took about twice as long. A 256-row chunk wants 32 to 128 threads per row, and the 1,024-thread version that was fastest for one row took 14.3 µs instead of 8.2, 1.7 times as long. At 4,096 rows every version moves the same 25 MB through memory and they all finish within 10% of each other. Reading four numbers at a time beat two in most cells of the table.

Two more things show up in the table. The fastest version at small shapes changes from run to run among versions within a few percent of each other, so a tuner’s pick among near-ties is partly noise; only the large gaps are stable. At these small shapes the calls sit near a floor of about 4 µs, likely the cost of issuing a call, which is why so many tie. And this is one kernel with two knobs: a matmul has tile sizes, pipelining depth and split strategies, and an engine build tunes every layer, which is why builds were slow.

Shape-specific tuning is also why an engine had a shape profile. A version tuned for the profile’s shapes could be the slowest choice outside them. TensorRT-LLM’s PyTorch runtime keeps the idea without the build step: its enable_autotuner option, on by default, runs a forward pass at startup “to populate the autotuner cache”, timing each tunable operation’s candidate kernels (“tactics”) for buckets of input shapes, and its documentation warns that “performance may significantly degrade if set to false”.

Batching and the KV cache

The rest of TensorRT-LLM’s runtime has the same shape as vLLM’s. In-flight batching is its name for continuous batching, covered in vLLM Architecture: requests join and leave the running batch at every step instead of waiting for a whole batch to finish. The KV cache is paged, as in Paged Attention, with 32 tokens to a block by default (tokens_per_block, “a power of two greater than 1”). As of v1.2.1:

• Block reuse is on by default. Blocks are stored “in a radix search tree as soon as they are filled”, so a new request reuses any cached prefix, the structure SGLang and RadixAttention introduced. Partial reuse is on by default too: when only some of a block’s tokens match, “the matched tokens are copied into the new block”, so a match can end at any token. Blocks can be offloaded to CPU memory, and evicted by priority.
• Chunked prefill is off by default. When it is on, a long prompt is computed in chunks between decoding steps, and a chunk must be a multiple of the block size, so that no block is left half-written by one chunk.
• The parallelism is fixed at launch: tensor, pipeline and expert parallelism are set when the server starts, not per request.

Quantization

Precision was the third step of a build: a quantized format and its scales, fixed in the engine. The formats are now produced by NVIDIA’s Model Optimizer (ModelOpt), a separate library whose README describes deployment “in … SGLang, TensorRT-LLM, TensorRT, or vLLM”, and TensorRT-LLM runs the pre-quantized checkpoints it writes. Quantization from Scratch derives the scale and zero point and works through GPTQ and AWQ; what matters here is which formats run on which GPUs:

This is the documentation’s support matrix as of October 2026. INT8 SmoothQuant, which earlier versions supported, is no longer in it.

Speculative decoding

Speculative Decoding, Derived explains the method: a cheap step proposes several tokens, one pass of the full model checks them all, and the accepted ones are exactly what the model would have produced. TensorRT-LLM’s list of proposers has changed with the backends. In the stable v1.2.1, Medusa and lookahead decoding run only on the engine backend (their configurations reject pytorch), and both are gone from the latest pre-release, v1.3.0rc29. What the PyTorch runtime supports in v1.2.1: EAGLE-3, multi-token prediction (MTP, the draft heads DeepSeek models are trained with), n-gram matching against the prompt and a separate draft model; the 1.3 pre-releases add PARD, DFlash and a suffix automaton.

Leaving the engine behind

The move took three steps. v0.17.0 (February 2025) added an experimental PyTorch backend. v1.0.0 (September 2025) made it the default, in the release notes’ words “Promote PyTorch to be the default LLM backend”. And v1.3.0rc21 (July 2026), a pre-release, removed the engine backend’s Python modules (“Remove Python modules and tests for the legacy TensorRT backend”), after deprecation warnings since rc11; in rc29, asking for it fails with “Unknown backend … Supported backends are ‘pytorch’”. The current stable release, v1.2.1 (April 2026), still ships trtllm-build. The project’s migration guide states the new arrangement: “PyTorch is now the sole execution backend”, with no engine-build step; in its words, “HuggingFace checkpoints load directly”.

The release notes do not give a reason, but the cost of the old design is visible in its own workflow. Every model architecture had to be rewritten in TensorRT-LLM’s network API, and every new method, a new attention variant or speculative proposer, had to be built for the engine as well, while vLLM and SGLang supported new models from their PyTorch code. The PyTorch runtime removes that second implementation.

What remains of compiling ahead is what the measurements above isolate:

• Fused modules: the gated MLP and add-and-normalize kernels, and custom attention, GEMM and mixture-of-experts kernels, written once and used from PyTorch.
• Auto-tuning at startup, per shape bucket, with a cache that can be saved to a file and reloaded.
• CUDA graphs, on by default for decoding, captured at startup for a list of batch sizes: a small, per-process version of a shape profile.

torch.compile support exists but is off by default and marked “prototype”. The project’s model engine runs all of this in one warmup, “calling specialized warmup methods for torch.compile, the autotuner, and CUDA graphs”. The cost the build step paid once per GPU and shape profile is now paid once per process, at startup.

Serving and NVIDIA’s numbers

trtllm-serve is an OpenAI-compatible server. It supports disaggregated serving, in which prompts and decoding run on separate GPU pools and the KV cache moves between them (over NIXL); the motivation is in Serving Metrics and Prefill–Decode Disaggregation. NVIDIA’s Dynamo can also run it, for multi-node deployments. The Triton Inference Server backend for engines went with the engine backend in rc21. And TensorRT-LLM still runs only on NVIDIA GPUs, on Linux.

Few independent, like-for-like comparisons with other engines exist, and none for current versions on current GPUs; the headline numbers are NVIDIA’s. In MLPerf Inference v5.1 (September 2025), NVIDIA’s Blackwell Ultra GB300 NVL72 system, serving DeepSeek-R1 with TensorRT-LLM, reported 5,842 tokens per second per GPU offline and 2,907 in the server scenario, which NVIDIA’s blog put at 45% and 25% above its GB200 NVL72 submission. These numbers measure a whole system, hardware and software, tuned by its vendor; they say what the stack can reach, not how TensorRT-LLM compares with vLLM or SGLang on the same GPUs.

Limits

• The measurements are a stand-in. TensorRT-LLM cannot run on a Mac, so the fusion and tuning above use Metal kernels on an Apple M4, and the scripts run only on Apple silicon. The mechanisms are the same; the sizes of the effects on an NVIDIA GPU are not, and on a fast GPU fusion matters more for decoding than here.
• Two fusions of many. An engine also fused attention, the QKV projections and quantization steps; the 26% fewer operations is for these two alone.
• One kernel tuned. Real tuning spans matmuls with many more choices.

Where this leads

• vLLM vs. SGLang vs. TensorRT-LLM compares the three engines, now all PyTorch runtimes.
• Running LLMs on a Laptop follows llama.cpp, which runs on the kind of machine these measurements came from.

Takeaways

• An engine build did three things: fused operations into fewer kernels, timed candidate kernels at the shapes in its profile and kept the fastest, and fixed a quantized format. The engine was tied to one GPU model, TensorRT version and shape profile.
• Fusing Qwen2.5-1.5B’s RMSNorms and MLP activations cut operations per decoded token from 1,425 to 1,055 and prefill time by 7%. Decoding improved only 4% on a laptop, where reading 3.09 GB of weights per token sets the pace; on a fast GPU, fusion matters more.
• The fastest version of a kernel depends on the shape: the version fastest for one row took 1.7 times as long as the best at 256 rows, and at small shapes the worst version took twice as long as the best.
• Since v1.0 TensorRT-LLM runs on PyTorch by default; the 1.3 pre-releases since rc21 (July 2026) remove the engine backend, while the stable v1.2.1 keeps it. Fusion, auto-tuning and CUDA graphs survive, done at startup instead of in a build.
• The KV cache is paged, 32 tokens to a block, with block reuse through a radix tree, with partial reuse, on by default. Medusa and lookahead decoding went with the engine backend.

References

Start with the migration guide for what changed, then the developer overview for how the PyTorch runtime is put together.

• NVIDIA, TensorRT Backend Removal: the migration guide; what was removed and what replaces it.
• NVIDIA, TensorRT-LLM Developer Guide: Overview: the PyTorch runtime’s executor, scheduler and KV cache manager.
• NVIDIA, TensorRT-LLM Release Notes and GitHub releases: the dates of each step; the pre-releases are only on GitHub.
• NVIDIA, KV Cache System: block reuse, the radix search tree, offloading and priority eviction.
• NVIDIA, Speculative Decoding: the proposers the PyTorch runtime supports.
• NVIDIA, TensorRT Quick Start Guide: what the TensorRT builder does with a network.
• NVIDIA, Model Optimizer: the quantization library behind the formats above.
• NVIDIA, NVIDIA Blackwell Ultra Sets New Inference Records in MLPerf Debut (2025): the MLPerf v5.1 DeepSeek-R1 numbers.
• NVIDIA, H100 Tensor Core GPU: the 3.35 TB/s memory bandwidth used above.