canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP

canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP

W4A16 INT4 routed experts + FP8 block 128×128 attention + BF16 Multi-Token Prediction (MTP) draft head retained — the first DeepSeek-V4-Flash quantization that ships a working MTP block, giving ~1.5× speculative decoding (spec-decode) speedup at bs=1 with no quality cost. Extends the W4A16-FP8 predecessor by patching the transformers calibration path so the MTP block survives the load.

TL;DR

Family / related artifacts

Why this exists

The W4A16-FP8 predecessor and `RedHatAI/DeepSeek-V4-Flash-NVFP4-FP8` both drop the MTP block because transformers 5.8.1's DeepseekV4PreTrainedModel declares:

_keys_to_ignore_on_load_unexpected = [r"(^|\.)mtp\..*"]

which silently filters every mtp.* tensor at from_pretrained time — without warning, without error. Calibration pipelines that go through from_pretrained produce quantized main weights paired with an absent MTP block; serving falls back to plain decode, losing the ~1.5–2× spec-decode speedup that V4-Flash's architecture provides.

This artifact bypasses the silent drop, runs the full 8-rank GPTQ calibration on a 768-sample corpus against the main routed experts, preserves the MTP block unquantized in BF16, and produces a serving artifact where speculative decoding actually fires.

Architecture & precision

Base model

Component precisions

Hardware validated

Same artifact, no weight changes between SKUs. Both validated cuda graphs ON.

Benchmarks

All numbers from the same artifact, vLLM HEAD 50d9dd902 + 4 patches cherry-picked (PRs #43248 / #43288 / #43290 / #43319).

Quality

Sampling: greedy, temperature 0. Methodology disclosed per row.

¹ The prior 30.0% AIME number was an lm-eval-harness aime24 task artifact — completions-mode prompt (no chat template), exact_match scorer on a thinking-mode model whose answers are wrapped in <think>…</think> + \boxed{N}. The scorer matched the literal answer string and missed virtually every correct response. A 1-shot smoke under proper chat-templated thinking=high methodology returned 2024-II-4: pred=33, exp=33, correct in 2072 completion tokens — model behavior is correct. Full 30-problem re-bench attempted 2026-05-24 on this RTX PRO 6000 box hit a reproducible CUDA illegal memory access (Worker_TP2: torch.AcceleratorError) under any concurrent thinking-mode load (cuda graphs and --enforce-eager both crash; concurrency=4 dies at ~11 min, concurrency=8 dies at ~90 s). Single-shot inference works. Re-bench deferred to H200 with jasl/vllm@ds4-sm120-experimental@abad5dc71 (the build the original Card D H200 numbers used).

² IFEval re-bench attempted 2026-05-24 hit the same RTX PRO 6000 stability issue. Deferred alongside AIME.

Throughput

vllm bench serve random 256-in / 256-out, MTP-spec num_speculative_tokens=1 (k=1 cap on this build — see Honest limitations), cuda graphs ON.

Per-replica, RTX PRO 6000 wins output throughput at every batch size; H200 still wins per-token TPOT median.

MTP draft-token acceptance per workload

Same artifact, bs=1, k=1.

Spec-decode wins at low concurrency (single-user interactive). At bs≥4 the verifier is already filling its batch lane, so extra verifier passes add overhead without saving wall-clock — matches the sibling artifact's framing of bs=1 as the headline operating point.

Cost per output token (node-level)

Boxes priced for cloud-rented hardware. Single-replica numbers measured; multi-replica totals are linear extrapolation.

At bs=1 (interactive), RTX PRO 6000 2×TP=2 is ~2.7× cheaper than H200 4×TP=2. At bs=16 the gap narrows because H200's per-replica throughput scales better with batch — H200 wins absolute throughput when you can fill it; RTX wins on $/token unless you genuinely need >1500 tok/s aggregate output.

Cross-validation: 2026-05-29 fresh Docker on RTX PRO 6000 ✓

Hardware coverage: all numbers below are from RTX PRO 6000 Blackwell Server Edition (SM 12.0a) on a Brev g7e.24xlarge. The image is expected to work on Workstation Edition (same SM 12.0a, same Marlin native cubins, same model + serve path) but we have not directly verified it ourselves. Reference TP=2 Workstation numbers from jasl's bench harness (`baselines/20260512_sm120_deployment_1c20f1a6d`) confirm the underlying stack runs on Workstation. Expect a 5-15% throughput delta from clock/memory-bandwidth differences. If anything misbehaves on Workstation Edition, open an issue at the repo.

Full bench matrix on canada-quant/dsv4-w4a16-rtxpro6000:v1 (the HF-published Docker image, built from jasl/vllm@27fd665b + canada-quant BF16-MTP cherry-pick + Marlin MoE c_tmp/workspace patches). All AIME runs at max_tokens = max_model_len - 500 = 65036 so reasoning runs to natural stop:

Headlines from this run:

• —Zero CUDA illegal-memory-access in 240 AIME thinking-mode problems across c=4 chat/high/max + c=1 high on both TP=2 and TP=4 = the Marlin MoE concurrent-decode race is fixed by the c_tmp clamp removal in PR vllm#43730 (which is baked into the v3 image via jasl/vllm@27fd665b).
• —TP=4 is 7-12× faster than TP=2 at AIME (chat 53m→5m, high 152m→13m, max 177m→26m). MoE expert sharding across 4 GPUs decisively wins.
• —Thinking-max regresses correctness AND triples wall (TP=4: high 29/30 in 13m vs max 27/30 in 26m). The artifact's sweet spot is reasoning_effort=high.
• —MTP holds 91-93% across all thinking modes and TP configs — the BF16-retained draft head is doing its job everywhere.

Raw JSON + per-bench logs in the reproduction repo.

Tuning attempts that DID NOT win on TP=4 Server (documenting so you don't repeat them)

We A/B-tested adopting jasl's TP=2 Workstation env tunings at TP=4 Server — none of them transferred. Stick with the v3 image defaults:

The image's defaults are the optimal config for TP=4 RTX PRO 6000 Server Edition as of 2026-05-29. If you're deploying on TP=2 Workstation Edition, jasl's reference config (`sm120_tp2_serve.env.example`) is the right starting point — it was tuned on that exact hardware.

Quick start

RTX PRO 6000 Blackwell — Docker (recommended)

The pre-built `canada-quant/dsv4-w4a16-rtxpro6000:v1` image bakes the full 13-layer recipe (jasl/vllm@27fd665b + canada-quant BF16 MTP cherry-pick + Marlin MoE c_tmp/workspace patches + cute.arch.fmin shim). ~3-5 min from docker load to a working endpoint on a g7e.24xlarge.

# 1. Pull the image tarball (~14 GB compressed)
hf download canada-quant/dsv4-flash-w4a16-rtxpro6000-image \
    --include "*.tar.gz" --local-dir .
docker load < dsv4-w4a16-rtxpro6000-v1.tar.gz

# 2. Cache the W4A16 model onto NVMe (~159 GB, ~1-2 min via xet on Brev)
HF_HOME=/opt/dlami/nvme/hf-cache hf download \
    canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP

# 3. Pull the serve helper
git clone https://github.com/canada-quant/dsv4-flash-w4a16-fp8-mtp.git
cd dsv4-flash-w4a16-fp8-mtp

# 4. Serve TP=2 (or TP=4 with --gpus all -e TP=4 -e MAX_NUM_SEQS=16)
docker run -d --gpus '"device=0,1"' --name dsv4-w4a16-serve \
    --shm-size=16g --ipc=host -p 8000:8000 \
    -v /opt/dlami/nvme/hf-cache:/root/.cache/huggingface \
    -v $(pwd)/scripts:/workspace/scripts:ro \
    -e TP=2 -e MAX_NUM_SEQS=4 -e MAX_MODEL_LEN=65536 -e GPU_MEM_UTIL=0.95 \
    canada-quant/dsv4-w4a16-rtxpro6000:v1 \
    bash /workspace/scripts/serve_rtx6000pro_w4a16.sh

# 5. Wait for /v1/models (~3-5 min model load + cudagraph capture)
until curl -sf http://127.0.0.1:8000/v1/models >/dev/null; do sleep 5; done

# 6. Run the full bench matrix (AIME chat/high/max + GSM8K + throughput)
docker exec dsv4-w4a16-serve bash -c \
    "TAG=tp2_64k MAX_MODEL_LEN=65536 bash /workspace/scripts/bench_matrix.sh"

RTX PRO 6000 Blackwell — from-source install (advanced)

# 1. Bootstrap vLLM (~25 min for source build)
git clone https://github.com/canada-quant/dsv4-flash-w4a16-fp8-mtp.git
cd dsv4-flash-w4a16-fp8-mtp
bash scripts/bootstrap_rtx6000pro.sh

# 2. Extra pins
source ~/venv-serve/bin/activate
pip install --quiet "flashinfer-python==0.6.8.post1" "flashinfer-cubin==0.6.8.post1" \
    "numba==0.65.0" "tilelang==0.1.9" "apache-tvm-ffi==0.1.9" "fastsafetensors>=0.2.2"

# 3. Apply patches
python scripts/patch_v4_forcausal_packed_mapping.py "$(python -c 'import vllm; print(vllm.__path__[0])')"
python scripts/patch_mtp_packed_mapping.py        "$(python -c 'import vllm; print(vllm.__path__[0])')"
python scripts/patch_nvidia_attn_scale.py         "$(python -c 'import vllm; print(vllm.__path__[0])')"
bash   scripts/patch_wo_a_bf16_path.sh             "$(python -c 'import vllm; print(vllm.__path__[0])')"

# 4. Download artifact (159 GiB) — already dequant'd in-artifact as of 2026-05-24,
#    no local preprocessing step required.
hf download canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP \
    --local-dir /scratch/weights/w4a16-fp8-mtp-gptq

# 5. Serve TP=2 (or TP=4 with 0,1,2,3)
CUDA_VISIBLE_DEVICES=0,1 bash scripts/serve_rtx6000pro.sh \
    /scratch/weights/w4a16-fp8-mtp-gptq 8000 2

Required runtime env vars on SM 12.x (already set inside `serve_rtx6000pro.sh` but worth knowing):

export VLLM_TRITON_MLA_SPARSE=1
export VLLM_TRITON_MLA_SPARSE_HEAD_BLOCK_SIZE=4
export VLLM_USE_FLASHINFER_SAMPLER=0

Without VLLM_TRITON_MLA_SPARSE_HEAD_BLOCK_SIZE=4 the sparse-MLA Triton kernel can crash during warmup with RuntimeError: Triton Error [CUDA]: an illegal memory access in _dequantize_and_gather_k_kernel. The FlashInfer sampler is also broken on TORCH_CUDA_ARCH_LIST=12.0a — fall back to PyTorch-native via VLLM_USE_FLASHINFER_SAMPLER=0.

H200

vllm serve canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP \
    --tensor-parallel-size 2 \
    --kv-cache-dtype fp8 --block-size 256 \
    --max-model-len 4096 \
    --gpu-memory-utilization 0.80 \
    --no-enable-prefix-caching \
    --tokenizer-mode deepseek_v4 \
    --tool-call-parser deepseek_v4 --enable-auto-tool-choice \
    --reasoning-parser deepseek_v4 \
    --speculative-config '{"method":"mtp","num_speculative_tokens":1}' \
    --trust-remote-code

Quantization recipe

Calibration recipe identical to the W4A16-FP8 predecessor with one change: the modeling class is patched to remove mtp.* from _keys_to_ignore_on_load_unexpected before from_pretrained, so the MTP block survives the load and is written back to the artifact at BF16.

vLLM build

Common patches (all platforms)

RTX PRO 6000 Blackwell (SM 12.0) only

(Previously this list also required a compressor/indexer FP8 → BF16 dequant preprocess step run against the local artifact. As of 2026-05-24 the dequant is baked into the published artifact — see Changes.)

H200 deployments need only the four common patches.

Honest limitations

1. 1.k=1 cap on spec-decode — current vLLM build limits num_speculative_tokens to 1 due to DeepGemm kernel assertion next_n == 1 or next_n == 2 in smxx_fp8_fp4_paged_mqa_logits.hpp:233. vLLM passes next_n = num_speculative_tokens + 1, so practical k is 1. The FLASHINFER_MLA_SPARSE attention backend hits the same kernel-side assertion. With the assertion relaxed, expect bs=1 speedup to rise from 1.49× to ~1.85× (matching sibling NVFP4 artifact's k=2 published number).
2. 2.Concurrent thinking-mode workloads on RTX PRO 6000 produce token-corrupted output — under concurrency ≥ 2 with thinking=high (long-decode workloads like AIME), the Marlin W4A16 MoE decode kernel on SM 12.0 produces token-stream corruption (CJK / Cyrillic / garbled ASCII spliced into the model's reasoning trace). The same hardware + same vLLM build serving the NVFP4 sibling (`canada-quant/DeepSeek-V4-Flash-NVFP4-FP8-MTP` via flashinfer_trtllm MoE) is essentially clean on the same workload (1/30 vs 14/30 corrupted at c=4 thinking). The bug is specific to the W4A16 + Marlin MoE decode path on SM 12.0. Investigation isolated through 7 controlled tests (sparse-MLA topk-chunk size, MTP-off, matmuldecode-off, eager-mode, concurrency sweep, NVFP4 vs W4A16 path comparison). **Workaround on RTX PRO 6000:** for batched thinking-mode workloads, serve the NVFP4 sibling artifact instead. For sequential (c=1) thinking-mode or any batched chat-mode (no thinking), this W4A16-MTP artifact works cleanly (GSM8K-20 chat-mode sequential = 20/20 = 100%, MTP draft acceptance 92.46%). Full debug log + reproducible benches: [`docs/findings/sm12xtokencorruption20260524.md](https://github.com/canada-quant/dsv4-flash-w4a16-fp8-mtp/blob/main/docs/findings/sm12x_token_corruption_2026_05_24.md). Filed upstream as [jasl/vllm#12`](https://github.com/jasl/vllm/issues/12).
3. 3.toolcall15 -2 pts vs predecessor — model-routing regressions on chain-completion (TC-07 stopped mid-chain to ask a clarifying question) and multi-tool extraction (TC-06 returned both translations as content text instead of routing two translate calls). Quality-wise the model completes the underlying intent; the harness scores tool-call-protocol fidelity, not task completion. Not a parser issue (confirmed by replay through --tool-call-parser deepseek_v4).
4. 4.GSM8K -1.3 pts vs predecessor's 8-shot strict-match — within one SE, but technically below. Likely calibration-set sensitivity rather than recipe drift (recipe is identical, hardware differs).
5. 5.NVFP4 native kernels on RTX PRO 6000 not auto-selected — even though csrc/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu exists in upstream vLLM, the backend selector doesn't pick it (`vllm-project/vllm#31085`). Until that lands, the sibling NVFP4 artifact on this hardware would route through Marlin too. This artifact's W4A16 path is the tested choice for RTX PRO 6000.

Reproduction

Full pipeline at `canada-quant/dsv4-flash-w4a16-fp8-mtp`. From a fresh 8× H200 box:

# Phase 0 — bootstrap (venv-calib + venv-serve + vendor + apply patches)
bash scripts/bootstrap_p5en_h200.sh

# Phase 1 — download upstream + dequant to BF16-MTP source (~30 min, ~660 GB)
bash scripts/phase1_dequant.sh

# Phase 2 — GPTQ calibration (8 ranks, ~15h wall)
bash scripts/run_phase2.sh

# Phase 3 — postprocess (rename + config patch + FP32 restore + MTP aliases)
bash scripts/postprocess_phase2.sh

# Phase 4 — verify
python scripts/verify_option_y.py /scratch/weights/w4a16-fp8-mtp-gptq

# Phase 5 — serve (see Quick start above for serve command)

Upstream contributions filed during this work

Changes

Files in the artifact

• —4 sharded model-*.safetensors files + model.safetensors.index.json (159 GB total)
• —config.json — vLLM-compatible quantization_config with MTP block excluded
• —tokenizer.json, tokenizer_config.json, generation_config.json, chat_template.jinja — upstream DSV4-Flash
• —recipe.yaml — the llm-compressor GPTQ recipe
• —README.md — this file

Citation

@misc{canada-quant-dsv4-flash-w4a16-fp8-mtp-2026,
  title  = {DeepSeek-V4-Flash W4A16-FP8 with BF16 MTP retained for vLLM speculative decoding},
  author = {Canada Quant},
  year   = {2026},
  publisher = {Hugging Face},
  url    = {https://huggingface.co/canada-quant/DeepSeek-V4-Flash-W4A16-FP8-MTP}
}

License

MIT, inherited from upstream deepseek-ai/DeepSeek-V4-Flash. Review at the upstream repo before commercial deployment.

Acknowledgments

• —DeepSeek for the base model + MTP architecture + inference reference.
• —jasl (`jasl/vllm` and `jasl/vllm-ds4-sm120-harness`) for the vLLM build pins (ds4-sm120-experimental for H200; ds4-sm120-preview-dev for RTX PRO 6000 SM 12.0) and the benchmark harness.
• —[`canada-quant/DeepSeek-V4-Flash-W4A16-FP8`](https://huggingface.co/canada-quant/DeepSeek-V4-Flash-W4A16-FP8) (predecessor) for the proven recipe topology this artifact extends with MTP.
• —[`canada-quant/DeepSeek-V4-Flash-NVFP4-FP8-MTP`](https://huggingface.co/canada-quant/DeepSeek-V4-Flash-NVFP4-FP8-MTP) (sibling) for the alias-injection pattern and MTP acceptance methodology.
• —vLLM, llm-compressor, compressed-tensors, FlashInfer maintainers.