DeepSeek-V4.1 kernels and 1M-context engine optimizations for NVIDIA Ampere (SM80 / SM86) on vLLM V1.
DeepInfer-SM80 ports the core ideas of DeepSeek's DeepInfer, which targets Blackwell (tcgen05, TMEM, CuteDSL), to Ampere GPUs such as the A100, A800, RTX 3090 and RTX A6000. It ships as two vLLM patches plus the equivalent drop-in Python/Triton/TileLang modules.
Fixed in v0.1.1 (
e2b585f): every version from the initial release (2026-10-20) up to that commit can let some layers use another index source's top-k selection on contexts longer than about 512 tokens. A before/after A/B on 8× GA100 64GB (TP4 × PP2,DSpark K=5) could not tell the two apart in long-context copy fidelity or speed, but the fix restores DeepSeek's reference behaviour, so please update. See docs/KNOWN_ISSUES.md.
Measured on NVIDIA GA100 64GB (SM80, 74 SMs) nodes serving DeepSeek-V4.1-Flash with DSpark K=5 speculative decoding, before the e2b585f fix; the before/after A/B showed no measurable change. Kernel speedups are lower on 108-SM A100/A800 parts; full tables, charts and independent runs are in docs/BENCHMARKS.md.
6 GPUs, TP2 × PP3: baseline vLLM vs. DeepInfer-SM80
| Metric | Baseline vLLM | DeepInfer-SM80 | Change |
|---|---|---|---|
| 16K cold prefill | 3,349 tok/s | 4,037 tok/s | +20.5% |
| 64K cold prefill | 3,627 tok/s | 4,507 tok/s | +24.2% |
| 64K time-to-first-token | 13.2 s | 10.7 s | −18.9% |
| 1-stream decode (code) | 151.8 tok/s | 160.8 tok/s | +5.9% |
| GPU-resident KV capacity | 4.65M tokens | 4.65M tokens | no VRAM overhead |
8 GPUs, TP4 × PP2: 13.3M GPU-resident KV tokens (12.69 full 1M-token contexts), 153–180 tok/s single-stream code decode, 620 tok/s aggregate at 12 streams, ~3,460 tok/s cold prefill.
Kernel microbenchmarks (isolated, verified against the baseline kernels):
| Kernel | Speedup | Accuracy |
|---|---|---|
_combine_topk_swa_indices_kernel (T=1536) |
22.17× | bit-exact |
dequantize_and_gather_k_cache (16K / 48K context) |
4.02× / 7.64× | bit-exact |
mhc_fused_post_pre_delayed (T=12) |
2.89× | < 1.9e-3 |
_sparse_attn_decode_kernel (B=6..24) |
2.08–2.14× | bit-exact |
_sparse_attn_prefill_kernel (T=1536) |
1.17× | < 4.8e-3 |
The patches target vLLM commit 79a7108d9a (PR #56201) and must be applied in order.
# 1. Check out the exact vLLM base (only reachable through the PR ref)
git init vllm && cd vllm
git remote add origin https://github.com/vllm-project/vllm.git
git fetch --filter=blob:none origin refs/pull/56201/head
git checkout 79a7108d9aea27ddab99ce1779290d300b17fc23
cd ..
# 2. Apply both patches (validates the base commit and the result)
bash Deepinfer-SM80/scripts/apply_patches.sh --vllm-dir vllm
# 3. Build for Ampere only, then check the environment
cd vllm && TORCH_CUDA_ARCH_LIST="8.0 8.6" pip install -e . --no-build-isolation && cd ..
python Deepinfer-SM80/scripts/verify_env.py
# 4. Run the kernel microbenchmarks on one GPU
bash Deepinfer-SM80/scripts/run_microbench_single_gpu.shAlready running an SM80 vLLM tree? Copy the 30 drop-in modules from vllm_sm80_deepinfer/ over your vllm package and confirm with scripts/check_dropin_sync.py. See docs/REPRODUCTION.md for the full procedure, the Docker recipe and known pitfalls.
| Compute capability | GPUs | Notes |
|---|---|---|
SM80 |
A100 40GB / 80GB, GA100 64GB, A800 40GB / 80GB | 164 KB shared memory per SM |
SM86 |
RTX 3090 / 3090 Ti, RTX A6000, A40, A10 | 100 KB shared memory per SM; all kernels use ≤ 48 KB per CTA |
| Document | Contents |
|---|---|
| KNOWN_ISSUES.md | Correctness issues, who is affected, how to update |
| CHANGELOG.md | Notable changes by date |
| REPRODUCTION.md | Base commit and fork lineage, patching, building, verification, pitfalls |
| BENCHMARKS.md | End-to-end serving, kernel microbenchmarks, multi-GPU topology comparison |
| ARCHITECTURE.md | Kernel-by-kernel design, forward-pass diagram, repository layout |
| 01 YOCO prefill compaction | Decoder SWA bounded compaction on TP2 × PP3 |
| 02 DSpark MTP | Variable-length verification, confidence head, top-k drafting |
| 03 DeepSelect attention | Sparse MLA, indexer and KV compression kernels |
| 04 mHC / Engram / MoE | Lazy mHC fusion, Engram prefetch, Marlin MoE, zero-overhead scheduling |
| 05 Topology & 1M context | TP2 × PP3 vs PP6 vs TP4 × PP2 capacity and latency analysis |
Built on DeepInfer (DeepSeek), vLLM and the haosdent/vllm SM80 enablement, TileLang, Triton, FlashInfer and Marlin. To cite this work, use CITATION.cff.
Licensed under the Apache License, Version 2.0.