[Tracking] DeepGEMM SM 12.x kernel coverage gaps for DeepSeek-V4-Flash on consumer Blackwell (RTX 50 / GB10)

[tonyliu312]

SM 12.x kernel coverage gaps for DeepSeek-V4-Flash on consumer Blackwell (RTX 50 / GB10)

Filed against jasl/DeepGEMM and tracked alongside vLLM #40899.

This is a comprehensive map of what's missing to run DeepSeek-V4-Flash on SM 12.x end-to-end through DeepGEMM, gathered from a 10-layer end-to-end debugging pass on dual NVIDIA GB10 (DGX Spark, SM 121, aarch64) using jasl/DeepGEMM@7a7a41a1 + jasl/vllm:ds4-sm120@8d0ebb76 against vLLM nightly 0.19.2rc1.dev220+g7b1bc0a3e.

What works in jasl/DeepGEMM today

The fork already includes SM 120-native CuTeDSL kernels for these V4 paths:

• sm120_tf32_hc_prenorm_gemm (mHC pre-norm GEMM — V4's hyper-connection layer)
• sm120_fp8_einsum
• sm120_fp8_paged_mqa_logits
• MmaMXF4NVF4Op (with the cute_dsl/warp/mma.py sm_120a whitelist patch — NOT in fork)

These paths boot end-to-end on RTX Pro 6000 (SM 120) per the PR description.

What's still missing for V4-Flash on SM 12.x

These dispatch sites in the fork still hard-code arch_major == 9 (Hopper) or arch_major == 10 (datacenter Blackwell) and reject SM 120/121, OR route to SM 10 kernels that use tcgen05.* instructions which SM 12.x silicon does not implement.

1. Dispatch routing — pure dispatch fixes (safe to extend to SM 12.x today)

File	Line	Site	Suggested fix
csrc/utils/layout.hpp	76	get_default_recipe: SM 10 → (1, 1, 128)	accept arch_major == 12 for the same recipe (verified working on GB10)
csrc/apis/layout.hpp	48, 56, 106, 110	SF transformation for FP4	accept arch_major == 12 (verified GB10)
csrc/jit/compiler.hpp	NVCCCompiler + NVRTCCompiler include_dirs	cutlass headers from third-party/cutlass/include	both compiler classes need -I third-party/cutlass/include added — currently NVCC fails with fatal error: cutlass/detail/helper_macros.hpp: No such file or directory

These alone unblock the next layer of failures.

2. The hard wall — kernel-source-level gaps (require actual SM 12.x kernel writing)

File	Site	Symptom on GB10
csrc/apis/gemm.hpp:99	fp8_fp4_gemm_nt dispatch only knows sm90_fp8_gemm_* and sm100_fp8_fp4_gemm_1d1d	Routing to sm100_fp8_fp4_gemm_1d1d produces ptxas: Instruction 'tcgen05.fence' not supported on .target 'sm_120f', Feature '.block32' not supported on .target 'sm_120f'. Needs sm120_fp8_fp4_gemm_1d1d kernel implementation — there is no SM 120 equivalent yet in the fork.
csrc/apis/attention.hpp:67, 177, 367	FP4 attention dispatch	Same tcgen05.* ISA mismatch as above. Needs SM 120 FP4 attention kernel(s).
csrc/apis/einsum.hpp:55	FP4 einsum dispatch	Same problem. Needs SM 120 FP4 einsum kernel (the existing sm120_fp8_einsum covers FP8 not FP4).

3. Companion vLLM-side gates that also need to relax

These are independent of DeepGEMM but block the deployment from reaching DeepGEMM in the first place. We've submitted the corresponding vLLM PRs:

vLLM file	What	PR
vllm/platforms/cuda.py support_deep_gemm	accept is_device_capability_family(120)	#41062
vllm/model_executor/layers/fused_moe/experts/deep_gemm_moe.py	DeepGemmFP4Experts._supports_current_device	#41062
vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py	_supports_current_device cap range < (11, 0) → < (13, 0)	#41028
vllm/.../csrc CMakeLists.txt MARLIN_*_ARCHS	add 12.0f family	#40923 (Approved by @Harry-Chen)

Reproduction command

# Inside vllm/vllm-openai:nightly-aarch64 with jasl/DeepGEMM@7a7a41a1 + the
# jasl/vllm ds4-sm120 branch synced into /usr/local/lib/python3.12/dist-packages/vllm/
vllm serve /models/DeepSeek-V4-Flash \
  --tensor-parallel-size 2 --distributed-executor-backend ray \
  --moe-backend deep_gemm \
  --kv-cache-dtype fp8_ds_mla \
  --gpu-memory-utilization 0.78 \
  --load-format instanttensor \
  --compilation-config '{"cudagraph_mode": "PIECEWISE"}' \
  --trust-remote-code

After applying section-1 patches, the tcgen05 failure surfaces during torch.compile / cudagraph capture (engine init, not at runtime).

Suggested rollup for cleaner re-submission

When you do the cleaner PR rebase you mentioned in #40969, please consider including:

1. Section 1 patches wholesale — they're literal or arch_major == 12 additions and the cutlass include for the JIT compiler (~10 lines total)
2. Authoring SM 120 FP4 kernels for the three Section-2 sites, or alternatively gating --moe-backend deep_gemm to fall through to the Triton sparse-MLA fallback on SM 12.x for FP4 ops until the kernels exist

Happy to test any rebased branch on dual GB10 TP=2 — quick turnaround (boot ≤ 2 min, TP=2 init ≤ 90s) and we already have the patch stack staged. cc @jasl @WoosukKwon @zyongye @lukealonso

— @tonyliu312 (filed from a debug trace on dual DGX Spark / SM 121 at 2026-04-28 01:25 Taipei)

[q77889966]

Just to add one more data point:
I tried jasl's vLLM branch on 8x RTX 5090, with tp=8.
It does run, but I only got around 8 tokens/s throughput in my setup.
So for RTX 5090, the current SM12x path seems to be runnable, but the performance is still extremely low / not really usable yet.

[hazyumps]

Field report from a debugging session on 2× NVIDIA GB10 (DGX Spark, SM 121, aarch64), DeepSeek-V4-Flash, TP=2 + --enable-expert-parallel over 200G ConnectX RoCE, MTP n=2, fp8 KV cache, vLLM v0.1.dev16581+gdda4668b5. Three findings in case they're useful for the rollup.

1. Deployment gotcha: NCCL silently falls back to TCP without RDMA device passthrough

On our container NCCL logged NET/IB : No device found → used NET/Socket for all cross-node collectives (all-reduce + EP all-to-all). libibverbs/libmlx5 were present, but the RDMA device nodes weren't mapped in (/dev/infiniband missing inside the container), so it ran every collective over the kernel TCP stack.

Fix was passing the devices into docker run:

--device=/dev/infiniband --cap-add=IPC_LOCK --ulimit memlock=-1:-1

→ NCCL switches to NET/IB (RoCE), single-stream throughput ~12 → ~34 tok/s. Non-obvious because the userspace libs being present makes it look configured. Might be worth a one-line note in the deployment docs.

2. sm12x_deep_gemm_fallbacks.py::_fp8_mqa_logits_topk_torch is FP32/cuBLAS-bound — a small bf16 tweak is a large win

With deep_gemm absent, the sparse-indexer top-k routes to this torch fallback, which dequants Q and K to FP32 and runs torch.matmul (cuBLAS SGEMM on CUDA cores, leaving Blackwell tensor cores idle). Under concurrency this dominated — 4 concurrent requests dropped to ~0.1 tok/s (effectively frozen), py-spy pinned in _fp8_mqa_logits_topk_torch → cublasSgemm.

Two no-kernel changes fixed it on our box:

• Dequant the matmul inputs to bf16 (tensor cores) and .float() the result — relu / accumulate / top-k stay fp32, so selection is unchanged (verified coherent output).
• Bump _SM120_MQA_LOGITS_MAX_SCORE_BYTES 64 MiB → 1 GiB — the 64 MiB cap forces many tiny chunk-loop iterations (launch overhead scales with KV length); GB10 has 128 GB unified memory to spare.

Result: 4 concurrent ~0.1 → ~60 tok/s, single-stream 29 → 34, output verified coherent. Happy to open a PR with the diff if it'd help.

3. Next gap: _forward_sparse_mla_prefill_triton wedges on heavy-concurrent + long-prompt

Same Python-double-loop-over-Triton-chunks structure as the indexer (for token_start … for index_start: accumulate_indexed_sparse_mla_attention_chunk). Under heavy concurrent + long-prompt load it hangs — requests don't abort on client disconnect, GPU stays pegged at running=N with no prompt/gen progress, py-spy pinned in _forward_sparse_mla_prefill_triton → _forward_prefill. Looks like the next SM 12.x kernel-coverage item after the indexer.

Glad to test rebased branches on dual GB10 TP=2 — quick turnaround.

[jasl]

Field report from a debugging session on 2× NVIDIA GB10 (DGX Spark, SM 121, aarch64), DeepSeek-V4-Flash, TP=2 + --enable-expert-parallel over 200G ConnectX RoCE, MTP n=2, fp8 KV cache, vLLM v0.1.dev16581+gdda4668b5. Three findings in case they're useful for the rollup.

1. Deployment gotcha: NCCL silently falls back to TCP without RDMA device passthrough

On our container NCCL logged NET/IB : No device found → used NET/Socket for all cross-node collectives (all-reduce + EP all-to-all). libibverbs/libmlx5 were present, but the RDMA device nodes weren't mapped in (/dev/infiniband missing inside the container), so it ran every collective over the kernel TCP stack.

Fix was passing the devices into docker run:

--device=/dev/infiniband --cap-add=IPC_LOCK --ulimit memlock=-1:-1

→ NCCL switches to NET/IB (RoCE), single-stream throughput ~12 → ~34 tok/s. Non-obvious because the userspace libs being present makes it look configured. Might be worth a one-line note in the deployment docs.

2. sm12x_deep_gemm_fallbacks.py::_fp8_mqa_logits_topk_torch is FP32/cuBLAS-bound — a small bf16 tweak is a large win

With deep_gemm absent, the sparse-indexer top-k routes to this torch fallback, which dequants Q and K to FP32 and runs torch.matmul (cuBLAS SGEMM on CUDA cores, leaving Blackwell tensor cores idle). Under concurrency this dominated — 4 concurrent requests dropped to ~0.1 tok/s (effectively frozen), py-spy pinned in _fp8_mqa_logits_topk_torch → cublasSgemm.

Two no-kernel changes fixed it on our box:

• Dequant the matmul inputs to bf16 (tensor cores) and .float() the result — relu / accumulate / top-k stay fp32, so selection is unchanged (verified coherent output).
• Bump _SM120_MQA_LOGITS_MAX_SCORE_BYTES 64 MiB → 1 GiB — the 64 MiB cap forces many tiny chunk-loop iterations (launch overhead scales with KV length); GB10 has 128 GB unified memory to spare.

Result: 4 concurrent ~0.1 → ~60 tok/s, single-stream 29 → 34, output verified coherent. Happy to open a PR with the diff if it'd help.

3. Next gap: _forward_sparse_mla_prefill_triton wedges on heavy-concurrent + long-prompt

Same Python-double-loop-over-Triton-chunks structure as the indexer (for token_start … for index_start: accumulate_indexed_sparse_mla_attention_chunk). Under heavy concurrent + long-prompt load it hangs — requests don't abort on client disconnect, GPU stays pegged at running=N with no prompt/gen progress, py-spy pinned in _forward_sparse_mla_prefill_triton → _forward_prefill. Looks like the next SM 12.x kernel-coverage item after the indexer.

Glad to test rebased branches on dual GB10 TP=2 — quick turnaround.

Tested, but seems to have no benefit in my RTX Pro 6000 environment.

BTW, I have a optimization notes that record every non-success experiments https://github.com/jasl/vllm-ds4-sm120-harness/blob/main/docs/sm120_optimization_notes.md

[hazyumps]

For anyone bringing up DeepSeek-V4-Flash on dual GB10 / DGX Spark (sm_121), I put together an end-to-end reproduction repo: https://github.com/hazyumps/deepseek-v4-flash-gb10

It sits on top of @jasl's fork (the SM12x enablement is his work) and captures the GB10-specific tuning + runbook to get from "crashes / wedges / ~12 tok/s" to a stable 384K-context, ~31 tok/s single-stream config across two Sparks (TP=2 + expert-parallel over RoCE, MTP n=2, fp8 KV). A few GB10-specific findings that may be useful here:

• NCCL 2.30.4 is the wedge fix. The shm_broadcast / "No available shared memory broadcast block found in 60 seconds" deadlocks under load go away on 2.30.4 — but PyTorch's bundled NCCL (2.28.x) shadows the system one, so it needs an explicit LD_PRELOAD of the 2.30.4 libnccl.so.2. (torch.cuda.nccl.version() still reports the compiled/bundled version; trust the NCCL banner in the logs.)
• RDMA must be made visible to the container (--device=/dev/infiniband + IPC_LOCK + memlock=-1), or NCCL silently runs every collective over TCP (~12 → ~30+ tok/s once IB engages, via NET/IB in the logs).
• Indexer top-k fallback. On sm_121 the lightning-indexer top-k has no native kernel and lands in the torch fallback. Two fixes helped: (1) bf16 matmul inputs instead of FP32 — the FP32/cuBLAS path froze concurrency (4-way 0.1 → ~60 tok/s); GB10-specific, @jasl saw no benefit on RTX Pro 6000. (2) Routing the per-chunk logits through the existing fused tf32 Triton MQA-logits kernel instead of the bf16 cuBLAS head-loop — ~+29% prefill at 9k, and it matches an fp32 reference exactly on the top-512 (i.e. more accurate than the bf16 fallback, not just faster). Patched fallback file + a correctness gate are in the repo.

Happy to fold any of this into #41834 if useful. Thanks @jasl / @tonyliu312 for the SM12x work.