diff --git a/recipes/deepseek-v4-1-flash-sm80-deepinfer-kernel-port-6x-a100-dspark-k5/recipe.yaml b/recipes/deepseek-v4-1-flash-sm80-deepinfer-kernel-port-6x-a100-dspark-k5/recipe.yaml index 23f86e475..e00265eca 100644 --- a/recipes/deepseek-v4-1-flash-sm80-deepinfer-kernel-port-6x-a100-dspark-k5/recipe.yaml +++ b/recipes/deepseek-v4-1-flash-sm80-deepinfer-kernel-port-6x-a100-dspark-k5/recipe.yaml @@ -2,12 +2,12 @@ slug: deepseek-v4-1-flash-sm80-deepinfer-kernel-port-6x-a100-dspark-k5 title: 'DeepSeek-V4.1-Flash SM80 DeepInfer kernel port 6x CMP 170HX DSpark K5' engine: name: vllm - version: 'V1 + Deepinfer-SM80 patch @ 2a828fc2122b' + version: 'V1 @ vllm PR #56201 base 79a7108d9a + haosdent/vllm dsv4-flash-a100 SM80 base + Deepinfer-SM80 patch @ 8895510cae8f' source_type: fork git_url: 'https://github.com/HalfVulpes/Deepinfer-SM80' - pinned_commit: 2a828fc2122b6fc59805c0c98b6aaf78e7da9b48 - install_note: 'Option A: git apply patches/vllm-sm80-deepinfer.patch from a vLLM V1 checkout root (all 6 files, 1,340 lines). Option B: bind-mount/copy vllm_sm80_deepinfer/vllm/* over site-packages/vllm/ (docker bind-mount read-only overlays work - the CMP-recipe pattern). Verify with the three in-repo microbenchmarks (microbench_sm80_kernels.py / _part2 / _part3) which print bit-exactness + latency tables; charts via generate_charts.py; full dataset in benchmarks/results/sm80_benchmark_summary.json. Requirements: SM80 (A100 40/64/80 GB, A800, GA100) or SM86 (RTX 3090/3090 Ti, RTX A6000, A40/A10/A30 - all kernels <=48 KB SRAM/CTA); vLLM V1 engine; Triton + TileLang JIT (startup cache-warm included). Base checkpoint: official deepseek-ai/DeepSeek-V4.1-Flash FP8.' -summary: "DeepSeek-V4.1-Flash served on six 64 GB SM80 GPUs (GA100 die, owner-classified 170HX) with HalfVulpes' Deepinfer-SM80 patch set: a 6-file vLLM V1 overlay porting DeepSeek's DeepInfer kernel architecture (SM100 tcgen05/TMEM) to Ampere mma.sync tensor cores - unified 512-dim SM80 sparse-MLA decode/prefill kernels with bitshift UE8M0 scale decoding, a bit-exact single-CTA Triton TopK+SWA combiner (22.17x), direct-store KV-cache gather (7.64x @48K), SM80-aligned TileLang mHC routing (2.89x), and YOCO L20-L39 ragged indexer-metadata reuse (19 redundant kernel chains eliminated per chunk). TP2 x PP3 (14/14/12) + DSpark K=5 + 1M context, 4.65M-token GPU-resident KV. Author-measured: 64K cold prefill 4,507 tok/s (+24.2% vs baseline vLLM SM80), 64K TTFT 10.7 s (-18.9%), c1 DSpark decode 160.8 tok/s code / 188.8 tok/s 16K+ long-gen, c4 170.5 aggregate. Microbenchmarks bit-exact (0 diff) or within BF16 noise. Topology guide: PP6 = 2.1x slower single-stream; TP2+TP4 asymmetric hurts Marlin tiles." + pinned_commit: 8895510cae8f8f1733401686e67a4d1e11e7e94b + install_note: 'Engine pin chain as stated by the repo at 8895510cae: base = vllm-project/vllm PR #56201 base commit 79a7108d9aea27ddab99ce1779290d300b17fc23 (v0.28.1rc0); SM80 base enablement = haosdent/vllm dsv4-flash-a100 (deepseek-v4-a100-support, commits 62195e9..c3046d1) rebased onto 79a7108d9a, shipped as patches/00-vllm-pr56201-sm80-base.patch (NEW at e375de2, +15,359 lines); DeepInfer-SM80 kernel + YOCO/PP optimizations via patches/vllm-sm80-deepinfer.patch (MODIFIED at 8895510: 30 files, 4,289 lines). Option A (reproducible full build from upstream): git clone vllm; git fetch origin 79a7108d9aea27ddab99ce1779290d300b17fc23; git checkout it; git apply patches/00-vllm-pr56201-sm80-base.patch; then git apply patches/vllm-sm80-deepinfer.patch. Option B (existing SM80 base tree, 11cf4a67f9 on top of 79a7108d9a): apply only patches/vllm-sm80-deepinfer.patch, or copy/bind-mount all 30 drop-in files: cp -r vllm_sm80_deepinfer/vllm/* site-packages/vllm/ (docker bind-mount read-only overlays work - the CMP-recipe pattern). Why vllm-sm80-deepinfer.patch does not apply to latest upstream vllm/main: upstream PR #56741 (cf1584f373) renamed vllm/models/deepseek_v4_1/ to deepseek_v41/, and upstream main ships only SM100 Blackwell + AMD ROCm backends for DeepSeek-V4.1-Flash. Verify with the three in-repo microbenchmarks (microbench_sm80_kernels.py / _part2 / _part3) which print bit-exactness + latency tables; charts via generate_charts.py; full dataset in benchmarks/results/sm80_benchmark_summary.json. Requirements: SM80 (A100 40/64/80 GB, A800, GA100) or SM86 (RTX 3090/3090 Ti, RTX A6000, A40/A10/A30 - all kernels <=48 KB SRAM/CTA); vLLM V1 engine; Triton + TileLang JIT (startup cache-warm included). Base checkpoint: official deepseek-ai/DeepSeek-V4.1-Flash FP8.' +summary: "DeepSeek-V4.1-Flash served on six 64 GB SM80 GPUs (GA100 die, owner-classified 170HX) with HalfVulpes' Deepinfer-SM80 patch set: a 30-file vLLM V1 overlay porting DeepSeek's DeepInfer kernel architecture (SM100 tcgen05/TMEM) to Ampere mma.sync tensor cores - unified 512-dim SM80 sparse-MLA decode/prefill kernels with bitshift UE8M0 scale decoding, a bit-exact single-CTA Triton TopK+SWA combiner (22.17x), direct-store KV-cache gather (7.64x @48K), SM80-aligned TileLang mHC routing (2.89x), and YOCO L20-L39 ragged indexer-metadata reuse (19 redundant kernel chains eliminated per chunk). TP2 x PP3 (14/14/12) + DSpark K=5 + 1M context, 4.65M-token GPU-resident KV. Author-measured: 64K cold prefill 4,507 tok/s (+24.2% vs baseline vLLM SM80), 64K TTFT 10.7 s (-18.9%), c1 DSpark decode 160.8 tok/s code / 188.8 tok/s 16K+ long-gen, c4 170.5 aggregate. Microbenchmarks bit-exact (0 diff) or within BF16 noise. Topology guide: PP6 = 2.1x slower single-stream; TP2+TP4 asymmetric hurts Marlin tiles." license: Apache-2.0 model: name: DeepSeek-V4.1-Flash @@ -45,12 +45,16 @@ kv_pool_tokens: 4648480 config: launch: - 'git clone https://github.com/HalfVulpes/Deepinfer-SM80 && cd Deepinfer-SM80' - - 'git apply patches/vllm-sm80-deepinfer.patch' - - '# or bind-mount vllm_sm80_deepinfer/vllm/* over site-packages/vllm/ as read-only overlays (CMP-recipe pattern)' - - '# serve DeepSeek-V4.1-Flash with TP2 x PP3 + DSpark K5 (author topology): see docs/05_TOPOLOGY_1M_CONTEXT_ANALYSIS_SM80.md' + - 'git clone https://github.com/vllm-project/vllm.git && cd vllm && git fetch origin 79a7108d9aea27ddab99ce1779290d300b17fc23 && git checkout 79a7108d9aea27ddab99ce1779290d300b17fc23' + - '# Option A: apply BOTH patches in order (00- base first, then DeepInfer-SM80):' + - 'git apply /path/to/Deepinfer-SM80/patches/00-vllm-pr56201-sm80-base.patch' + - 'git apply /path/to/Deepinfer-SM80/patches/vllm-sm80-deepinfer.patch' + - '# Option B: on an existing SM80 base tree (11cf4a67f9 on 79a7108d9a) apply only patches/vllm-sm80-deepinfer.patch, or bind-mount/copy all 30 drop-in files: cp -r vllm_sm80_deepinfer/vllm/* /path/to/site-packages/vllm/' + - '# verify: python3 benchmarks/microbench_sm80_kernels.py && python3 benchmarks/microbench_sm80_part2.py && python3 benchmarks/microbench_sm80_part3.py' + - '# serve DeepSeek-V4.1-Flash 6x SM80 64GB: TP2 x PP3 VLLM_PP_LAYER_PARTITION=14,14,12 + DSpark K=5 + max context 1,048,576 (README Quick Start + docs/05 section 1; repo ships no serve script)' runtime: env: "Author topology on 6x 64 GB: TP2 x PP3 partition 14/14/12 (VLLM_PP_LAYER_PARTITION=14,14,12 per benchmarks/results/sm80_benchmark_summary.json), DSpark K=5 speculative, max context 1,048,576, 4.65M GPU-resident KV. Alternative topologies measured in docs/05: PP6 (TP1xPP6) = 65.5-68.5 tok/s c1 (2.15x slower), TP2+TP4 asymmetric = ~55.5 tok/s unspeculated (moie_intermediate 2304/4=576 hurts Marlin 4.5x128 tiling). Deployment env/args: the patch is engine-internal; no special env vars documented in the README (DSpark handled by vLLM's speculative-config; the mHC/Triton kernels self-configure per SM count)." - notes: 'Drop-in modules touch: vllm/v1/attention/ops/rocm_aiter_mla_sparse.py (unified 512-dim SM80 sparse MLA), vllm/models/deepseek_v4_1/amd/rocm.py (single-CTA TopK+SWA combiner), common/ops/cache_utils.py (bitshift UE8M0 KV gather), model_executor/kernels/mhc/{tilelang,warmup}.py (SM80 mHC + JIT warm), models/deepseek_v4_1/nvidia/model.py (YOCO L20-L39 metadata reuse). Numerics: 100% bit-exact on 4 of the 5 kernel families (0.00e+00 diff); mHC fused kernels <1.9e-3 (BF16 MMA order) and sparse-attn prefill <4.8e-3 - documented per-kernel, no accuracy regression claimed. Author rig also tiers KV to pinned /dev/shm RAM and 512 GiB NVMe (docs/05) - rig infra, not a patch requirement.' + notes: 'Drop-in modules (30 files at 8895510cae) touch: vllm/models/deepseek_v4_1/pp_share.py (PPSharePlan cross-PP YOCO shared-KV/indexer-K/candidate replication - missing in 2a828fc212, the PP patch set was BROKEN without it), nvidia/{dspark,vl_model}.py (TP-sharded DSpark on PP; VL PP input_ids propagation), common/engram.py (exact-size cudaHostRegister + pthread pread Engram host shards), ampere/ampere_sparse.py (SM80 SWA metadata builder, no FlashMLA dependency), attention.py (per-head Triton indexer warmup); vllm/v1/attention/backends/mla/{indexer,streamed_workspace}.py (bounded 1M-context streamed indexer workspace ~135 MB vs 5.28 GB), v1/attention/backends/utils.py, v1/attention/ops/mqa_logits_candidates_triton.py (candidate-only indexer scoring for L24/28/32/36), v1/core/kv_cache_utils.py + v1/core/single_type_kv_cache_manager.py + v1/core/sched/{scheduler,output,async_scheduler}.py, v1/structured_output/utils.py, v1/worker/gpu/model_runner.py + v1/worker/gpu/cudagraph_utils.py, v1/worker/gpu/sample/batch_shard.py, v1/worker/gpu/spec_decode/rejection_sampler.py + v1/worker/gpu/spec_decode/dspark/utils.py; plus the initial-release files v1/attention/ops/rocm_aiter_mla_sparse.py (unified 512-dim SM80 sparse MLA), models/deepseek_v4_1/amd/rocm.py (single-CTA TopK+SWA combiner), common/ops/cache_utils.py (bitshift UE8M0 KV gather), model_executor/kernels/mhc/{tilelang,warmup}.py (SM80 mHC + JIT warm), models/deepseek_v4_1/nvidia/model.py (YOCO L20-L39 metadata reuse), model_executor/layers/sparse_attn_indexer.py + model_executor/model_loader/weight_utils.py + model_executor/models/qwen3_dspark.py. Numerics: 100% bit-exact on 4 of the 5 kernel families (0.00e+00 diff); mHC fused kernels <1.9e-3 (BF16 MMA order) and sparse-attn prefill <4.8e-3 - documented per-kernel, no accuracy regression claimed. Author rig also tiers KV to pinned /dev/shm RAM and 512 GiB NVMe (docs/05) - rig infra, not a patch requirement.' performance: claimed_by: submitter metrics: @@ -84,7 +88,7 @@ artifacts: kind: engine_image pointer_type: git_commit pointer: 'https://github.com/HalfVulpes/Deepinfer-SM80' - pin: 'HEAD 2a828fc2122b6fc59805c0c98b6aaf78e7da9b48 at ingest 2026-10-21 (verified: 6 files, 1,340-line patch, README+summary JSON match)' + pin: 'HEAD 8895510cae8f8f1733401686e67a4d1e11e7e94b at revision 2026-10-28 (verified: 30 drop-in files; patches/00-vllm-pr56201-sm80-base.patch new +15,359 lines + patches/vllm-sm80-deepinfer.patch modified 30 files/4,289 lines; README+summary JSON match)' - kind: model_weights pointer_type: hf_model