The Blackwell Ceiling: Deploying DeepSeek-V4-Flash with Prefill-Decode Disaggregation and Confronting the sm_120 Fallback Kernel Bottleneck

Introduction

In the high-stakes world of large language model deployment, few challenges are as unforgiving as bringing a cutting-edge model to a new hardware architecture. This article chronicles a pivotal segment of an opencode coding session where an AI assistant deployed DeepSeek-V4-Flash — a 284-billion-parameter Mixture-of-Experts model with FP4 quantization and Multi-Token Prediction (MTP) — on eight NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture) using the SGLang inference engine. The headline architectural deliverable — prefill-decode (PD) disaggregation — was achieved, with the prefill phase running on GPUs 0–3 (NUMA0) and the decode phase on GPUs 4–7 (NUMA1), connected by NIXL/UCX KV cache transfer. But the performance numbers told a sobering story: approximately 10 tokens per second at batch size 1, and roughly 28 tok/s at concurrency 16 after extensive optimization, against a user target of ~1000 tok/s.

This article synthesizes the full arc of the work: the initial deployment, the baseline benchmark that revealed a 40× performance gap, the systematic optimization campaign that exhausted every configuration lever, and the definitive diagnosis that identified the root cause — the sm_120 fallback kernels for sparse-MLA attention and MoE slot-GEMV, which run on CUDA cores rather than tensor cores, leaving the GPU's tensor-pipe utilization below 1%. The session ultimately concluded that configuration tuning and checkpoint changes can only yield incremental gains against this structural ceiling, and that reaching the throughput target would require either disabling MTP to restore batch capacity or building a custom split-K tensor-core sparse-attention kernel — a multi-week engineering effort akin to earlier custom kernel work on the Kimi K2.6 model.

The Deployment: From Reconnaissance to PD Disaggregation

The segment opened with a dramatic pivot. The assistant had just completed an intensive optimization campaign on the Kimi K2.6 model — building a native C/C++/CUDA speculative-decoding engine, writing custom sm_120 verify-attention kernels, achieving CUDA graph capture-safety, and delivering a verified 3–6× decode speedup over the Triton baseline. Yet the bottleneck had shifted to MoE expert imbalance at batch size 1, a structural limitation of tensor parallelism with small batches. The user issued a new directive: pivot to DeepSeek-V4-Flash with PD disaggregation [2].

The assistant's response was methodical. Rather than diving into implementation, it launched a parallel reconnaissance phase [4], gathering information from three independent sources simultaneously: the HuggingFace model card (to understand the model's architecture and quantization format), the SGLang PD disaggregation documentation (to learn the deployment requirements), and the target machine's live state (GPU topology, disk space, SGLang version, running services). This triage reflected a deep awareness that premature implementation could lead to wasted effort.

The reconnaissance revealed a critical compatibility gap: the model card for DeepSeek-V4-Flash was "100% vLLM — forks + many patches, zero SGLang mention" [5]. The model uses the DeepseekV4ForCausalLM architecture with sparse attention mechanisms, a lightning indexer, compressor, and hash-based MoE routing — a fundamentally new architecture that SGLang 0.5.11 almost certainly did not support. A GitHub issue confirmed that users attempting to load the model hit unsupported architecture errors. The assistant also discovered that the root filesystem was 98% full with only 22 GB free, while the model required 159 GB — necessitating the deletion of the 548 GB Kimi K2.6 checkpoint [5].

Despite these obstacles, the assistant pressed forward. It freed disk space, pulled a fresh SGLang main branch, and built a complete environment with all required dependencies: flashinfer 0.6.12, sglang-kernel 0.4.3, tilelang, NIXL, and CUDA 13.0 toolchain [12][13][14]. The editable install of SGLang required careful dependency management — resolving import shadowing from leftover site-packages directories [25][26] and debugging a shebang path issue in the launch script [21]. After these hurdles, the model loaded successfully on a single-node TP4 configuration and generated correct output: "The capital of France is Paris." [28][29][30].

The headline deliverable — prefill-decode disaggregation — was then orchestrated [36][37]. The assistant stopped the single-node server, verified that all eight GPUs were freed, installed the missing sglang-router package (a Rust-based component that coordinates traffic between prefill and decode servers), and launched the split deployment. The prefill server ran TP4 on GPUs 0–3 (NUMA0) with NIXL as the transfer backend, while the decode server ran TP4 on GPUs 4–7 (NUMA1). A router on port 8000 coordinated traffic, and KV cache was transferred between the groups via NIXL/UCX. The end-to-end pipeline was verified: requests sent to the router were correctly dispatched to the prefill server, and the KV cache was transferred to the decode server for generation [36].

The Baseline That Changed Everything

With the model serving correctly, the assistant ran a baseline benchmark [32]. The results were devastating:

• Concurrency=1: 10.22 output tokens per second, mean TPOT of 94.13 ms, mean TTFT of 753.55 ms
• Concurrency=8: ~24 tok/s aggregate (truncated due to timeout) The model card for DeepSeek-V4-Flash on vLLM reports roughly 107 tok/s on comparable hardware. The assistant was seeing roughly one-tenth of the expected performance. The user's target was ~1000 tok/s — a 40× gap. The assistant's diagnostic process was systematic [33][34]. It confirmed that CUDA graphs were active (cuda graph: True in the startup logs), ruling out graph-capture issues. It checked for straggler processes, verified the MoE runner backend (marlin, which correctly routes to the sm_120 Triton fallback path), and confirmed that memory settings were reasonable. The diagnosis was clear: this was not a misconfiguration. The bottleneck was structural — the sm_120 architecture lacks the fast-path kernels that make DeepSeek-V4-Flash performant. The root cause was architectural. The model uses MXFP4 quantized experts (4-bit weights), which on sm_100 (Blackwell B200) would route through DeepGEMM and NVIDIA's trtllm-gen compiled kernels — highly optimized tensor-core paths. On sm_120 (the RTX PRO 6000 Blackwell variant), those fast paths are arch-gated: the CUDA capability check is_sm100_supported() returns False, and SGLang falls back to Triton-based kernels that run on CUDA cores (SIMT) rather than tensor cores. The forced SGLANG_FP8_PAGED_MQA_LOGITS_TORCH path computes logits through pure PyTorch operations. The result is an order-of-magnitude throughput penalty baked into the software stack [33][34].

The Optimization Campaign: Exhausting Every Lever

Faced with a 40× performance gap, the assistant launched a systematic optimization campaign, testing every available configuration lever:

NCCL LL+Ring Tuning: The assistant experimented with different NCCL communication protocols (LL vs Ring) to optimize the all-reduce operations required for tensor parallelism. The impact was negligible — the bottleneck was compute, not communication.

CUDA Graphs: Already active from the initial deployment. The assistant confirmed that graph capture was working correctly but that the underlying kernels were simply too slow.

Tilelang Indexer Fusion: The assistant attempted to use tilelang's JIT compilation to fuse the sparse attention indexer operations. This failed on sm_120 — the tilelang compiler did not support the architecture, producing JIT-compile errors.

Non-Marlin MoE Backends: The assistant explored alternative MoE runner backends (e.g., triton, cutlass). These were invalid for FP4 experts — the MXFP4 format requires the Marlin path for correct dequantization.

Expert Parallelism (EP): The assistant attempted to split experts across GPUs to reduce the per-GPU MoE computation. This actually made performance worse due to PCIe all-to-all overhead — the cross-GPU communication for expert routing dominated the compute time.

MTP/EAGLE Speculative Decoding: The assistant enabled the model's native Multi-Token Prediction heads for speculative decoding. This delivered a 47% single-request throughput boost — significant, but not enough to close the gap. At higher concurrency, the benefit vanished because the MTP verifier saturated the GPU, halving the effective batch size [34].

NVFP4 Quantization: The most impactful change was switching from the W4A16 (INT4 weights, FP8 attention) checkpoint to an actual NVFP4 checkpoint. NVIDIA's NVFP4 format routes MoE execution through tensor-core paths (Marlin W4A16 or cutlass FP4 grouped GEMM), whereas the W4A16 path fell through to Triton fallback. NVFP4 delivered ~28 tok/s at C=16 — a 24% improvement over the baseline [34].

Despite these efforts, the ceiling remained. The assistant had exhausted every configuration lever available within the SGLang framework, and the best throughput achieved was ~28 tok/s — still roughly 40× below the user's target.

The Definitive Diagnosis: sm_120 Fallback Kernels

The assistant ran a definitive GPU profile to understand where the time was going. The profile traced 63% of decode time to a single kernel: _tiled_sparse_decode_kernel, the sm_120 Triton fallback for sparse MLA attention. This kernel launches only 64 blocks (1 batch × 64 heads) on approximately 170 streaming multiprocessors, serially iterating through all 512 top-k tokens. The result is catastrophic under-utilization: the majority of GPU compute units sit idle while a tiny fraction do the work [34].

The profile also revealed that the MoE slot-GEMV kernel consumed 39% of GPU time, also running on CUDA cores rather than tensor cores. The NVFP4 quantization switch had improved this somewhat by routing the grouped GEMM through tensor-core paths, but the sparse attention kernel remained the dominant bottleneck regardless of MoE backend.

The assistant's analysis identified two compounding factors:

1. Low occupancy: The sparse attention kernel launches only 64 blocks on ~170 SMs. With each block serially iterating 512 top-k tokens, the effective parallelism is far below what the hardware can sustain. A custom split-K kernel could increase block count by splitting the top-k dimension across more blocks, improving occupancy and utilizing the GPU's tensor cores.
2. MTP verifier overhead: The MTP speculative decoding heads, while providing a 47% single-request boost, halved the effective batch size at higher concurrency because the verifier consumed GPU resources that could otherwise serve additional requests. The assistant calculated that disabling MTP would increase batch capacity from C=8 to C=16, potentially doubling throughput — but this would also eliminate the speculative decoding benefit, resulting in a net wash. The decisive finding was that no amount of configuration tuning could close the gap. The sm_120 fallback kernels are not slow due to suboptimal parameters; they are slow because they run on the wrong compute units. The fast fused DSA/MoE stack (DeepGEMM, trtllm-gen, FP4 indexer) is arch-gated to SM100, and the sm_120 path is a compatibility layer maintained for correctness, not performance.

The Path Forward: Custom Kernel Engineering

The assistant concluded the segment with a clear-eyed assessment of the path forward. The structural ceiling imposed by sm_120 fallback kernels means that configuration tuning and checkpoint changes can only yield incremental gains. To reach the user's throughput target, one of two approaches would be required:

Approach 1: Disable MTP to restore batch capacity. By disabling the MTP speculative decoding heads, the verifier overhead is eliminated, allowing the system to serve larger batches. The assistant calculated that this could increase effective batch size from C=8 to C=16, potentially doubling throughput. However, this would also eliminate the 47% single-request speculative decoding benefit, resulting in roughly equivalent performance at low concurrency.

Approach 2: Build a custom split-K tensor-core sparse-attention kernel. This is the high-risk, high-reward path. The assistant had precedent for this approach from the earlier K2.6 work, where a custom verify-attention kernel achieved 3–6× decode speedup over Triton. A similar effort for DeepSeek-V4-Flash's sparse MLA attention would involve writing a custom CUDA kernel that:

• Uses split-K decomposition to increase block count from 64 to thousands
• Routes computation through tensor cores (sm_120a FP4/FP8 tensor instructions)
• Fuses the indexer gather with the attention matmul to reduce memory traffic
• Achieves capture-safety for CUDA graph integration The assistant estimated this as a multi-week engineering effort, comparable in scope to the K2.6 custom kernel work. The potential reward is substantial: if the attention kernel could be accelerated by 3–6× (as the K2.6 verify kernel was), and combined with NVFP4 tensor-core MoE, the system could potentially reach 100–200 tok/s — still short of the 1000 tok/s target, but a meaningful improvement.

Themes and Lessons

Several overarching themes emerge from this segment:

The sm_120 architecture gap. The NVIDIA RTX PRO 6000 (sm_120) is a Blackwell-generation GPU, but it is not a B200. The sm_120 architecture lacks certain tensor-core capabilities that the sm_100/sm_103 variants provide, and the software stack — SGLang, vLLM, DeepGEMM — has optimized primarily for the sm_100 path. The sm_120 path is a fallback, maintained for correctness but not performance. This creates a frustrating situation where the model runs correctly but at a fraction of the hardware's potential.

The discipline of measurement. The assistant's systematic approach to benchmarking — establishing a baseline before optimizing, testing each configuration lever independently, and running definitive GPU profiles — exemplifies best practices in performance engineering. The 10.22 tok/s baseline was disappointing, but it was invaluable as a reference point for measuring improvement.

The hard ceiling of configuration tuning. The optimization campaign tested NCCL settings, CUDA graphs, tilelang fusion, MoE backends, expert parallelism, speculative decoding, and quantization formats. Each delivered incremental gains, but none could close the fundamental gap. This is a case study in the limits of configuration-level optimization when the bottleneck is architectural.

The value of custom kernel engineering. The assistant's earlier work on K2.6 demonstrated that custom CUDA kernels can deliver 3–6× speedups over Triton fallbacks. The DeepSeek-V4-Flash deployment faced the same structural bottleneck, and the same solution — custom kernel engineering — was identified as the path forward. This reinforces the lesson that for novel hardware architectures, the off-the-shelf inference stack may not deliver production-grade performance.

Scope discipline. Perhaps the most impressive aspect of the session was the assistant's ability to accept the bottleneck and pivot to the achievable deliverable. Rather than spending days chasing kernel optimizations, the assistant recorded the baseline honestly, installed NIXL, and built the PD disaggregation architecture — the work that actually needed to get done. The sm_120 performance story was documented for future reference, but the project continued moving forward.

Conclusion

This segment of the opencode session tells a story of ambition meeting reality. The assistant successfully deployed DeepSeek-V4-Flash on Blackwell GPUs with prefill-decode disaggregation — a complex architectural achievement involving NUMA-aware process placement, NIXL/UCX KV cache transfer, and router-coordinated request dispatch. The model generated correct output, the CUDA graphs captured, and the disaggregated pipeline functioned end-to-end.

But the performance numbers revealed a hard truth: the sm_120 architecture's fallback kernels impose a structural ceiling that no amount of configuration tuning can overcome. The fast fused DSA/MoE stack that makes DeepSeek-V4-Flash performant on B200 hardware is arch-gated to SM100, and the Triton fallback kernels for sparse-MLA attention and MXFP4 MoE run on CUDA cores at a fraction of the hardware's potential.

The assistant's response — measure systematically, optimize exhaustively, diagnose definitively, and document honestly — is a model of disciplined engineering. The path forward is clear: either disable MTP to restore batch capacity, or invest in custom kernel engineering to build a split-K tensor-core sparse-attention kernel. Either way, the session has established a precise understanding of the bottleneck, a quantified baseline, and a roadmap for achieving the user's throughput targets. The sm_120 ceiling is real, but it is not insurmountable — it simply requires the right tools and the right scope of work.## References

[1] "The Architecture of a Status Report: How a Native DDTree Inference Engine Took Shape on Blackwell GPUs" — Article on the comprehensive status document that preceded the pivot to DeepSeek-V4-Flash.

[2] "The Pivot to DeepSeek-V4-Flash: A Strategic Experiment in Prefill-Decode Disaggregation" — Analysis of the user's directive to pivot from K2.6 to DeepSeek-V4-Flash.

[3] "The Pivot: From Custom DDTree Kernels to DeepSeek-V4-Flash PD Disaggregation" — The assistant's structured information-gathering response to the pivot directive.

[4] "The Intelligence-Gathering Phase: Parallel Reconnaissance Before a Complex Model Deployment" — Examination of the parallel reconnaissance operations launched to understand the model, deployment architecture, and target system.

[5] "The Architecture Pivot: How One Message Uncovered a Fundamental Compatibility Gap in Deploying DeepSeek-V4-Flash on SGLang" — The discovery that SGLang lacked support for the DeepSeek-V4 architecture.

[12] "The Moment the Model Arrives: A Pivotal Transition in Deploying DeepSeek-V4-Flash on Blackwell" — The model download and environment setup phase.

[13] "The Pivot Point: From Download to Deployment on Blackwell" — Transition from download to deployment.

[14] "Mapping the Unknown: The Reconnaissance Message That Defined DeepSeek-V4-Flash Deployment on Blackwell" — The reconnaissance that defined the deployment strategy.

[21] "The Shebang Trap: How a Hardcoded Python Path Nearly Derailed a DeepSeek Deployment" — Debugging a hardcoded Python path in the launch script.

[25] "The Ghost in the Package: Diagnosing a Shadowed Editable Install" — Diagnosing import shadowing from leftover site-packages.

[26] "The Phantom Directory: Debugging Python Import Shadowing in a Multi-GPU ML Deployment" — Further debugging of import shadowing.

[28] "The First Smoke Test: Launching DeepSeek-V4-Flash on Blackwell GPUs" — The first successful model launch.

[29] "The Waiting Game: Monitoring DeepSeek-V4-Flash's First Breath on Blackwell" — Monitoring the first startup.

[30] "The Moment of Truth: Validating DeepSeek-V4-Flash on Blackwell GPUs" — Validation of correct model output.

[32] "The Baseline That Changed Everything: Benchmarking DeepSeek-V4-Flash on Blackwell sm_120" — The baseline benchmark that revealed the 10× performance gap.

[33] "The Moment of Reckoning: Accepting the sm_120 Ceiling on DeepSeek-V4-Flash" — The diagnostic process that identified the sm_120 fallback kernel bottleneck.

[34] "The Honest Baseline: Diagnosing SM120 Performance Ceilings in the DeepSeek-V4-Flash Deployment" — The definitive diagnosis of the sm_120 ceiling.

[35] "The Gatekeeper: Validating NIXL for Prefill-Decode Disaggregation on Blackwell" — Verification of NIXL/UCX for KV cache transfer.

[36] "The Transitional Moment: Orchestrating Prefill-Decode Disaggregation for DeepSeek-V4-Flash on Blackwell" — The planning and orchestration of PD disaggregation.

[37] "The Transitional Pivot: Cleaning State and Installing the Router for PD Disaggregation" — Cleaning state and installing the missing router package.