DeepInfer-SM80: High-Performance DeepSeek-V4.1 Kernels & 1M-Context Engine Optimizations for NVIDIA Ampere (SM80 / SM86)

DeepInfer-SM80 ports and adapts core architectural concepts and kernel optimizations from DeepSeek's DeepInfer (SM100 / Blackwell tcgen05.mma + TMEM + CuteDSL) to NVIDIA Ampere (SM80 / SM86) GPUs—including NVIDIA A100 (40GB / 64GB / 80GB), A800, RTX 3090, and RTX A6000—running on the vLLM V1 engine.

While upstream DeepInfer relies on Blackwell-exclusive hardware features (tcgen05 warp-group instructions, Tensor Memory TMEM, hardware UE8M0 scale instructions, and 228 KB shared memory per SM), DeepInfer-SM80 re-engineers the critical execution path around Ampere's mma.sync.aligned.m16n8k16 Tensor Cores, 164 KB L1/SRAM capacity (SM80), IEEE-754 integer bitshift UE8M0 scale decoding, SM-count-aware grid scheduling, and YOCO (L20) shared-KV reuse.

Scaling from 6 GPUs (TP2 × PP3) to 8× NVIDIA A100/GA100 64GB (SM80) in a balanced 2-stage pipeline (TP4 × PP2, VLLM_PP_LAYER_PARTITION=20,20, --gpu-memory-utilization 0.95, max_num_seqs=12) places all 20 encoder layers (L0–L19) on PP0 and all 20 YOCO consumer layers (L20–L39 + 3 DSpark EAGLE draft layers) on PP1—completely eliminating cross-stage KV replication of Layer 20 (compress_ratio=1) and unlocking 13,306,380 GPU-resident KV tokens (12.69 × 1M full context windows, 2.86× capacity vs 6 GPUs):

• Long-Context Single & Multi-Stream Cold Prefill (TP4 × PP2):
  • 40.6K prompt (C=1): 3,541.5 tok/s (11.47 s TTFT)
  • 162.2K prompt (C=1): 3,436.1 tok/s (47.20 s TTFT)
  • 4 × 162.2K prompts (C=4, 649K total tokens): 3,471.2 tok/s aggregate (187.05 s wall span)
  • 649.0K single-stream prompt (C=1): 2,759.9 tok/s (235.16 s TTFT, 63.27 GiB / 64.00 GiB peak VRAM on PP1)

Validated on a 6× NVIDIA A100/GA100 64GB (SM80) node serving DeepSeek-V4.1-Flash (TP2 × PP3 + DSpark K=5 Multi-Token Speculative Decoding + 1,048,576 max context length, 4.65M GPU-resident KV tokens):

Every modified kernel was benchmarked in isolation and verified against baseline reference outputs on SM80 hardware:

![SM80 Kernel Speedups](charts/01_sm80_kernel_speedups.png)
<!-- slide -->
![SM80 Sparse MLA Decode Scaling](charts/02_sparse_mla_decode_scaling.png)
<!-- slide -->
![Per-Chunk Prefill Latency Breakdown](charts/03_prefill_chunk_breakdown.png)

flowchart TD
    subgraph Prefill["Chunked Prefill & Decode Forward Pass (40 Backbone + 3 DSpark EAGLE Layers)"]
        A["Input Hidden States + 4-Stream mHC Residual [T, 4, 5120]"] --> B["TileLang mHC Pre-Layer Sinkhorn Routing<br/>(Aligned 6×8×128 FMA + Tensor Core T.gemm @ T>=16)"]
        B --> C{"Layer Type?"}
        C -->|"Encoder L0–L19<br/>(compress_ratio = 4)"| D["Lightning Indexer + Triton Combine TopK+SWA<br/>(Single-CTA Disjoint Interval Store: 22.17x faster)"]
        C -->|"YOCO Producer L20<br/>(compress_ratio = 0)"| E["Build Ragged K Metadata Once + Full KV Write<br/>(Cached in ForwardContext for L21–L39)"]
        C -->|"YOCO Consumers L21–L39<br/>(shared_kv_source = L20)"| F["Reuse L20 Ragged K Metadata + Shared KV Read<br/>(Eliminates 19x Redundant Indexer Metadata Launches)"]
        D --> G["Unified 512-dim SM80 Sparse MLA Kernel<br/>(Bitshift UE8M0 + 2 CTAs/SM + Direct-Out Split-1)"]
        E --> G
        F --> G
        G --> H["MoE / MLP + TileLang mHC Fused Post+Pre Boundary<br/>(83 Boundaries/Step: 2.89x Faster at T=12)"]
    end

File: vllm/v1/attention/ops/rocm_aiter_mla_sparse.py

DeepSeek-V4.1's DeepSelect Sparse Multi-head Latent Attention (kv_lora_rank = 512, qk_rope_head_dim = 64, v_head_dim = 512) stores the 448-dim NoPE latent in float8_e4m3fn with UE8M0 power-of-two block scales (128-element groups) alongside a 64-dim bfloat16 RoPE suffix.

• Bitshift IEEE-754 UE8M0 Scale Decoding: Replaces transcendental tl.math.exp2(scale_u8.to(tl.float32) - 127.0) calls on the SFU (4 instructions/warp) with a single integer bitshift (scale_u8.to(tl.int32) << 23).to(tl.float32, bitcast=True) (1 instruction on the integer pipe), producing bit-identical IEEE-754 float32 values for all valid UE8M0 exponents.
• Unified Single-Pass 512-Dim Inner Loop: Instead of splitting 448-dim NoPE (7 × 64 iterations) and 64-dim RoPE into separate loads and dot products, our kernel loads the full 512-dim tile ([0:448] FP8 dequantized in registers via tl.where(d < 448, ...) + [448:512] BF16 RoPE) and executes a single tl.dot(q, k) Tensor Core MMA.
• 2 CTAs/SM Occupancy (num_warps=4) + Direct-Out Fast Path (num_kv_splits=1): Tuning BLOCK_N=64, num_warps=4, num_stages=2 halves per-CTA register pressure so 2 CTAs co-reside per SM80 SM, hiding HBM gather latency across sparse Top-K indices. When num_kv_splits == 1, Stage 1 writes directly to the final output tensor and skips the Stage 2 reduction kernel entirely (2.14× decode speedup, 0.00e+00 bit-exact match).

File: vllm/models/deepseek_v4_1/amd/rocm.py

Previously, combine_topk_swa_indices executed a chain of PyTorch GPU tensor allocations (torch.arange, torch.clamp, torch.minimum, torch.full, torch.cat) on every sparse attention layer during prefill (1,083.7 µs at T=1536).

• We replaced the entire chain with a single Triton kernel (_combine_topk_swa_indices_kernel, 1 CTA per query token, num_warps=4).
• Critical Multi-Warp Correctness Guarantee: In Triton, tl.store vectors of different power-of-two widths (512, 128, 256) distribute the same linear column index to different warps. Overlapping an initial -1 padding store with subsequent out_topk and swa_vals stores without an explicit barrier creates an inter-warp write race whenever topk_len < TOPK_WIDTH. We partition the output row into three strictly disjoint column intervals:
  1. [0, topk_len) written exclusively by the compressed Top-K indices
  2. [topk_len, topk_len + swa_len) written exclusively by the SWA window indices
  3. [topk_len + swa_len, COMBINED_TOPK) padded with -1
• This achieves a 22.17× speedup (1,083.72 µs → 48.88 µs at T=1536, saving 31.0 ms per 1,536-token prefill chunk) with 100% bit-exact output across all sequence lengths and compression ratios (0, 1, 2).

File: vllm/models/deepseek_v4_1/common/ops/cache_utils.py

• Replaces exp2 scale decoding with bitshift UE8M0 (scale_raw.to(tl.int32) << 23).to(tl.float32, bitcast=True).
• Eliminates intermediate torch.empty + .copy_() buffer allocations by passing out_k_nope and out_k_rope strides directly into the Triton kernel, achieving a 4.02× speedup at 16K context and 7.64× speedup at 48K context.

Files: vllm/model_executor/kernels/mhc/tilelang.py, vllm/model_executor/kernels/mhc/warmup.py

DeepSeek-V4.1 executes mHC 4-stream residual routing (hc_mult = 4, hidden_size = 5120) at 83 boundaries per forward pass (2 × 40 backbone layers + 3 DSpark EAGLE layers).

• Aligned Split Geometry for hidden_size = 5120: The stock _SMALL_FMA_CONFIG = (6, 16, 128) splits 5120 into 16 slices of 320 elements (320 % 128 = 64 != 0), failing alignment checks and falling back to an under-occupied 10-block launch on SM80 (13.5% of an A100's SMs!). Adding (6, 8, 128) (5120 // 8 = 640 = 5 × 128, launching 40 blocks) and lowering _SMALL_M_THRESHOLD from 32 → 16 so T >= 16 dispatches to Tensor Core T.gemm delivers a 2.89× speedup at T=12 (357.7 µs → 123.8 µs) and 1.30× at T=6.

File: vllm/models/deepseek_v4_1/nvidia/model.py

• In DeepSeek-V4.1's You-Only-Cache-Once (YOCO) architecture, Layers 21–39 (shared_kv_source_layer = 20, compress_ratio = 1) all read from Layer 20's full-resolution KV cache and share identical seq_lens_expanded, block_table, and compress_ratio.
• Caching (k_ragged_offset, k_ragged_SlotMapping, total_k) in forward_context at L20 eliminates 19 redundant build_ragged_k_Metadata GPU kernel chains per prefill chunk.

Deepinfer-SM80/
├── README.md                          # Overview, benchmarks, architecture, and usage guide
├── LICENSE                            # Apache-2.0 License
├── CITATION.cff                       # Citation metadata (DeepInfer-SM80, DeepInfer, vLLM)
├── patches/
│   ├── 00-vllm-pr56201-sm80-base.patch # Base SM80 enablement patch on top of vllm PR #56201 (79a7108d9a)
│   └── vllm-sm80-deepinfer.patch      # DeepInfer-SM80 kernel & YOCO optimization patch (all 6 files, 1,529 lines)
├── vllm_sm80_deepinfer/               # Drop-in optimized vLLM Python/Triton/TileLang modules
│   └── vllm/
│       ├── model_executor/kernels/mhc/
│       │   ├── tilelang.py            # SM80-aligned TileLang mHC FMA & Tensor Core kernels
│       │   └── warmup.py              # Startup JIT cache warming for split=4,8,16 configs
│       ├── models/deepseek_v4_1/
│       │   ├── amd/rocm.py            # Single-CTA bit-exact _combine_topk_swa_indices_kernel
│       │   ├── common/ops/cache_utils.py # Bitshift UE8M0 + direct-store KV gather kernels
│       │   └── nvidia/model.py        # YOCO L20–L39 ragged indexer metadata reuse
│       └── v1/attention/ops/
│           └── rocm_aiter_mla_sparse.py  # Unified 512-dim SM80 Sparse MLA decode/prefill kernels
├── benchmarks/
│   ├── microbench_sm80_kernels.py     # Isolated SM80 microbenchmark: Sparse MLA & KV gather
│   ├── microbench_sm80_part2.py       # Isolated SM80 microbenchmark: TileLang mHC kernels
│   ├── microbench_sm80_part3.py       # Isolated SM80 microbenchmark: TopK+SWA combiner
│   ├── generate_charts.py             # Reproducible 300-DPI Matplotlib chart generator
│   └── results/
│       └── sm80_benchmark_summary.json # Complete kernel & E2E benchmark dataset
├── charts/
│   ├── 01_sm80_kernel_speedups.png
│   ├── 02_sparse_mla_decode_scaling.png
│   ├── 03_prefill_chunk_breakdown.png
│   ├── 04_e2e_production_throughput.png
│   └── 05_8gpu_tp4pp2_scaling_1_2_4_12.png
└── docs/
    ├── 01_YOCO_PREFILL_COMPACTION_SM80.md
    ├── 02_DSPARK_MTP_VARLEN_TOPK_SAMPLING_SM80.md
    ├── 03_DEEPSELECT_ATTENTION_KV_KERNELS_SM80.md
    ├── 04_MHC_ENGRAM_MOE_ZOS_SM80.md
    └── 05_TOPOLOGY_1M_CONTEXT_ANALYSIS_SM80.md

1. Upstream Package Rename (PR #56741): Upstream vllm-project/vllm PR #56741 (commit cf1584f373) renamed vllm/models/deepseek_v4_1/ to vllm/models/deepseek_v41/.
2. SM80 Base Enablement: Upstream vllm/main currently ships only the Blackwell (SM100 FlashInfer/CuteDSL) and AMD ROCm (gfx942/gfx950) backends for DeepSeek-V4.1-Flash. Running DeepSeek-V4.1-Flash on NVIDIA Ampere (SM80/SM86) requires the SM80 base enablement layer derived from:
   • Upstream vLLM Base: vllm-project/vllm PR #56201 ([Model] Support DeepSeek-V4.1-Flash) at commit 79a7108d9aea27ddab99ce1779290d300b17fc23 (v0.28.1rc0 base).
   • Community SM80 Fork: haosdent/vllm (dsv4-flash-a100 / deepseek-v4-a100-support, commits 62195e9..c3046d1) rebased onto 79a7108d9aea27ddab99ce1779290d300b17fc23 + V4.1 SM80 Sparse MLA / software-FP8 / DSpark-on-PP enablement (provided in full in patches/00-vllm-pr56201-sm80-base.patch).

Clone vllm, check out the exact PR #56201 commit (79a7108d9aea27ddab99ce1779290d300b17fc23), and apply the two patches in order:

git clone https://github.com/vllm-project/vllm.git
cd vllm
git fetch origin 79a7108d9aea27ddab99ce1779290d300b17fc23
git checkout 79a7108d9aea27ddab99ce1779290d300b17fc23

# Step 1: Apply the SM80 base enablement layer (haosdent/vllm dsv4-flash-a100 rebased onto PR #56201)
git apply /path/to/Deepinfer-SM80/patches/00-vllm-pr56201-sm80-base.patch

# Step 2: Apply the DeepInfer-SM80 kernel & YOCO optimizations
git apply /path/to/Deepinfer-SM80/patches/vllm-sm80-deepinfer.patch

If you are already running the SM80 base branch (11cf4a67f9 on top of 79a7108d9a):

git apply /path/to/Deepinfer-SM80/patches/vllm-sm80-deepinfer.patch

Or copy/bind-mount the 6 drop-in files directly into your vllm Python package directory:

cp -r vllm_sm80_deepinfer/vllm/* /path/to/site-packages/vllm/

To verify bit-exactness and measure kernel latency on your own A100 / RTX 3090 / A6000 GPU:

python3 benchmarks/microbench_sm80_kernels.py
python3 benchmarks/microbench_sm80_part2.py
python3 benchmarks/microbench_sm80_part3.py
python3 benchmarks/generate_charts.py

When deploying DeepSeek-V4.1 (40 Backbone + 3 DSpark EAGLE layers) across 6× or 8× 64GB/80GB Ampere GPUs for simultaneous low-latency single-user experience and multi-stream 1M context, see our full analysis in docs/05_TOPOLOGY_1M_CONTEXT_ANALYSIS_SM80.md:

This work builds directly upon the architectural innovations and open-source foundations of the following projects:

1. DeepSeek Team — DeepInfer: Original DeepSeek-V4.1 inference engine, DeepSelect Sparse MLA architecture, YOCO shared-KV design, DSpark multi-token speculative decoding, and mHC (Manifold-Constrained Hyper-Connections) algorithms.
2. vLLM Project & haosdent/vllm (dsv4-flash-a100): PagedAttention, vLLM V1 engine scheduler, DeepSeek-V4.1 PR #56201 (79a7108d9a), and initial SM80 A100 enablement.
3. TileLang & OpenAI Triton: DSL compilers enabling high-performance SM80 Tensor Core and warp-level kernel synthesis in Python.
4. FlashInfer & IST-DASLab Marlin: Sparse MLA and MXFP4/W4A16 Ampere Tensor Core GEMM primitives.

@software{deepinfer_sm80_2026,
  author       = {HalfVulpes},
  title        = {{DeepInfer-SM80: High-Performance DeepSeek-V4.1 Inference Kernels and 1M-Context Engine Optimizations for NVIDIA Ampere (SM80/SM86)}},
  year         = {2026},
  url          = {https://github.com/HalfVulpes/Deepinfer-SM80},
  license      = {Apache-2.0}
}

Licensed under the Apache License, Version 2.0.