Why 72GB Blackwell GPUs Can't Run DeepSeek V4: The SM120 Kernel Gap
Forum topic · 小凯 · 2026-05-26
Summary
A practitioner trying to load DeepSeek-V4-Flash on two RTX Pro 5000 cards (72GB GDDR7 each, 144GB combined) hits a RuntimeError "Unsupported architecture" despite sufficient VRAM. The root cause is not memory but a missing kernel path: DeepGEMM ships SM90 and SM100 implementations but no SM120 versions, and Blackwell consumer/workstation GPUs use compute capability 12.0 (SM120), which differs from the data-center SM100 in Tensor Core instructions such as tcgen05.fence and wgmma. Two operators are missing—tf32_hc_prenorm_gemm for DeepSeek's manifold-constrained hyper-connections and fp8_paged_mqa_logits for the Lightning Indexer—so DeepSeek V4 cannot fall back to a generic CUDA or Triton path. Workarounds include SGLang with TileLang (partial), community flags that produce all-zero tokens, and Marlin W4A16 quantization at roughly 5 tok/s instead of the expected 50+. The article compares alternative hardware choices and notes that upstream fixes have no ETA as of May 2026.
Key points
- Hardware has enough memory, but software does not run. Two RTX Pro 5000 cards provide 144GB of GDDR7, which is theoretically sufficient for DeepSeek-V4-Flash (~148GB) under TP=2. The model still fails to load with
RuntimeError: Assertion error: Unsupported architecture.
- The actual cause is an SM120 kernel gap. NVIDIA Blackwell splits into two lines: data-center parts use SM100 (compute capability 10.0), while consumer and workstation parts—RTX 5090, RTX Pro 5000, RTX Pro 6000—use SM120 (12.0). Tensor Core instructions such as
tcgen05.fence (SM100) and wgmma (SM90) are not present on SM120, and SM120 requires its own warp-matrix-multiply-accumulate rewrite.
- Two DeepGEMM operators are missing for SM120. The DeepGEMM library contains
sm90_* and sm100_* paths for both required kernels but no sm120_* files:
tf32_hc_prenorm_gemm — used by DeepSeek V4's Manifold-Constrained Hyper-Connections (mHC) pre-norm layer.
fp8_paged_mqa_logits — used by the Lightning Indexer for FP8 paged multi-query attention logits.
No generic CUDA fallback and no Triton path exist; both kernels only live inside DeepGEMM.
- Workarounds and their limits.
- *SGLang + TileLang*: lets the mHC path JIT-compile for SM120, but
paged_mqa_logits has no replacement backend and still blocks inference.
- *ktransformers community patch*: reportedly boots on RTX Pro 4000 (24GB, SM120) with flags such as
SGLANG_DISABLE_DEEP_GEMM=1 and SGLANG_OPT_USE_TILELANG_INDEXER=1, but produces all-zero tokens or NaN logits due to bugs in the FP8 fallback (including an incorrect F.relu that clips negative attention scores).
- *Marlin W4A16 fallback*: runs, but only at about 5 tok/s on a 397B-class model versus 50+ tok/s for a proper FP4 path, and the activation distribution mismatch breaks speculative-decoding draft heads.
- *Waiting for upstream fixes*: open issues include DeepGEMM #236 (SM120 feature request, Feb 2026), DeepGEMM #317 (DeepSeek-V4 crash, Apr 2026), CUTLASS #3096 (SM120 TMA WS grouped GEMM failure), and a FlashInfer tracker needing 12 patches before SM120 compiles. No ETA as of 2026-05-26.
- Practical guidance for GPU selection.
- For production DeepSeek V4 deployment, use SM90 (H100/H200) or SM100 (B200) where vLLM and SGLang work out of the box.
- For local development, RTX Pro 6000 (96GB) gives more headroom than Pro 5000 (72GB) and is often the difference between fitting a model or not.
- SM120 hardware is powerful but its software ecosystem lags; it is a reasonable choice only for models without exotic kernel requirements or for users willing to act as early adopters.
- Historical pattern. The same gap appeared for SM89 (RTX 4090) at launch: community patches typically catch up within 3–6 months. Until then, users either accept fallback performance or move to SM90/SM100 cards.
References
- DeepGEMM Issue #317: DeepSeek-V4 on SM120 — Unsupported architecture
- DeepGEMM Issue #236: Feature Request: Support sm_120 (5090 and Blackwell 6000 Pro)
- ktransformers Issue #2001: RTX PRO 4000 Blackwell all-zero tokens
- CUTLASS Issue #3096: SM120 TMA WS grouped GEMM failure
- FlashInfer Issue #2577: NVFP4 mm_fp4 GEMM broken on SM120
This page is an English static mirror generated for search and AI citation.
It may be a full translation or structured summary of the Chinese original.
Canonical interactive discussion lives on the Chinese page:
https://zhichai.net/topic/177620818