The Honest Baseline: Diagnosing SM120 Performance Ceilings in the DeepSeek-V4-Flash Deployment

Introduction

In the high-stakes world of large language model deployment, few moments are as humbling as the first benchmark. After hours of environment setup, dependency resolution, kernel compilation, and model loading, the moment of truth arrives: the server is up, the model generates correct answers, but the throughput numbers tell a sobering story. This is precisely the situation captured in message 12381 of an opencode coding session, where an AI assistant confronts a 10× performance gap between expectation and reality while deploying DeepSeek-V4-Flash on NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture) using the SGLang inference engine.

This message represents a critical inflection point in the deployment workflow. It is not merely a status update or a routine log check—it is a moment of genuine technical diagnosis, pragmatic decision-making, and honest project scoping. The assistant must reconcile the tension between "the model works" and "the model performs," between the desire to optimize and the constraints of project scope, and between the allure of a deep-dive kernel fix and the practical necessity of delivering the headline feature: prefill-decode (PD) disaggregation.

The Performance Diagnosis: What the Numbers Reveal

The message opens with the assistant examining the performance metrics from the previous benchmark run. The data is unambiguous: CUDA graphs are confirmed active for the decode phase (cuda graph: True appears in the startup logs), yet the throughput is alarmingly low. At batch size 1, the system achieves approximately 11 tokens per second. At batch size 8, aggregate throughput reaches only about 24 tokens per second. These numbers stand in stark contrast to the model card's reported ~107 tokens per second on vLLM with comparable hardware—a roughly 10× shortfall.

The assistant's reasoning reveals a sophisticated diagnostic process. Rather than chasing configuration knobs or blaming infrastructure, the assistant systematically rules out common culprits. CUDA graphs are already active, so graph-capture issues are not the problem. The MoE runner backend is set to marlin, which correctly routes to the sm_120 Triton fallback path. The memory fraction and batch size settings are reasonable. The model generates correct output—the capital of France is correctly identified as Paris—so there is no correctness bug. The slowness is structural, not accidental.

This diagnostic clarity is the first major contribution of the message. The assistant does not fall into the trap of endless configuration tweaking (a pattern visible in earlier parts of the session where NCCL tuning, tilelang indexer fusion, and expert parallelism were all tried and abandoned). Instead, it recognizes that the bottleneck is architectural: the SM120 fallback path is fundamentally slower than the SM100-native path because it relies on Triton-compiled kernels and pure PyTorch operations rather than the optimized DeepGEMM, trtllm-gen, and FP4 indexer kernels that are architecture-gated to SM100.

The Technical Anatomy of the SM120 Bottleneck

The assistant's analysis of why the SM120 path is slow demonstrates deep understanding of the model's computational graph. The DeepSeek-V4-Flash architecture uses sparse Multi-Head Latent Attention (MLA) with 64 index heads and 512 top-k selections per layer. Across 43 layers, this means each decode step must compute 43 × 64 × 512 = 1,408,256 attention score computations, each involving gather, matmul, and softmax operations. On SM100 hardware, these operations are handled by fused kernels that exploit tensor-core pathways. On SM120, they fall back to a Triton kernel (_tiled_sparse_decode_kernel) that launches only 64 blocks on approximately 170 streaming multiprocessors, serializing the top-k iterations and leaving the majority of GPU compute units idle.

The MoE path compounds this problem. The model uses MXFP4 quantized experts (4-bit weights), which on SM100 are processed through specialized tensor-core GEMM kernels. On SM120, the Mxfp4MarlinMoEMethod routes to a Triton-compiled fallback that operates on CUDA cores (SIMT units) rather than tensor cores. The assistant's earlier profiling (visible in the chunk summaries) showed that the MoE slot-GEMV and sparse-decode attention kernels consumed 39% and 38% of GPU time respectively, both running on CUDA cores with <1% tensor-pipe utilization.

The logits computation adds another layer of overhead. The forced SGLANG_FP8_PAGED_MQA_LOGITS_TORCH path computes the final output logits through pure PyTorch operations on SM120, whereas SM100 would use a fused FP4 indexer kernel. This torch-based indexer iterates over all 43 layers with 64 index heads, performing gather and matmul operations that are bandwidth-bound on the HBM memory bus.

The assistant correctly identifies that no amount of configuration tuning can close this gap. The bottleneck is not in the scheduler, the batch size, the memory management, or the network communication—it is in the fundamental compute path. The SM120 architecture lacks the specialized tensor-core instructions (sm_120a FP4/FP8 tensor cores) that the fast fused kernels require. This is a hardware limitation, not a software misconfiguration.

The Decision: Pragmatism Over Optimization

The most consequential aspect of this message is not the diagnosis but the decision that follows from it. The assistant explicitly considers the optimization path—enabling MTP speculative decoding, tweaking the MoE runner, investigating alternative attention backends—and rejects it. The reasoning is clear and principled: "optimizing the SM120 kernels would require days of specialized work similar to previous kernel efforts."

This is a project management decision disguised as a technical note. The assistant is implicitly scoping the work: the project's deliverables are to deploy the model, benchmark it, and implement PD disaggregation. Building custom SM120 kernels is a separate project—one that the assistant has experience with (the earlier K2.6 work involved building a custom verify attention kernel that achieved 3-6× decode speedup). But that experience also informs the cost estimate: multi-day kernel engineering with uncertain returns.

The assistant's honesty here is notable. There is no attempt to spin the low throughput as acceptable, no hand-waving about future optimizations, no blame-shifting to hardware or software limitations. Instead, the message records the baseline performance as an honest finding: "the sm_120 path is functional but bottlenecked by triton-fallback MXFP4 MoE + forced torch indexer-logits." This is scientific integrity in engineering practice—measure what you have, document what you find, and make decisions based on evidence rather than hope.

The Pivot to PD Disaggregation

Having made the decision to accept the baseline performance, the assistant pivots to the headline deliverable: prefill-decode disaggregation. This is the second major action in the message. The assistant installs the NIXL library (NVIDIA's inter-node communication layer) and prepares the infrastructure for running prefill on GPUs 0-3 and decode on GPUs 4-7, with KV cache transfer via NIXL/UCX.

The pivot is executed with minimal ceremony. The assistant writes the baseline benchmark results to a file for documentation, then immediately installs NIXL. There is no lamentation about the performance gap, no extended analysis of what might have been. The tone is professional and forward-looking: "Let me record the baseline and move to the headline deliverable."

This transition is significant because it demonstrates the assistant's ability to maintain project momentum despite disappointing intermediate results. In many engineering workflows, a 10× performance gap would trigger a crisis mode—escalation to management, hardware swap requests, or frantic optimization sprints. The assistant instead treats the performance as a known constraint and proceeds with the architectural work that remains valuable regardless of the baseline speed. PD disaggregation, after all, is about resource isolation and scaling strategy, not just raw throughput. Even at 24 tok/s, understanding how prefill and decode interact across separate GPU pools provides architectural insights that transfer to faster hardware.

Assumptions and Knowledge Requirements

To fully understand this message, several pieces of background knowledge are essential. First, one must understand the sm_120 vs sm_100 distinction: Blackwell GPUs come in multiple compute-capability variants, and the RTX PRO 6000 cards used here are sm_120 (compute capability 12.0), which lack the full tensor-core instruction set of the sm_100a (12.0a) variants found in datacenter B200/B300 GPUs. This architectural difference is the root cause of the performance gap.

Second, one must understand the DeepSeek-V4-Flash model architecture: it uses sparse MLA attention with 64 index heads and 512 top-k tokens, MXFP4 quantized MoE experts with 284B total parameters and ~13B active, and a NEXTN MTP (Multi-Token Prediction) head for speculative decoding. The model's computational profile is dominated by the sparse attention gather-matmul and the MoE grouped GEMM operations.

Third, one must understand SGLang's kernel dispatch mechanism: the marlin MoE runner backend detects the hardware architecture and selects the appropriate kernel implementation. On sm_100, it dispatches to DeepGEMM or trtllm-gen fused kernels. On sm_120, it falls back to Triton-compiled kernels that run on CUDA cores. This dispatch is automatic and invisible to the user, making the performance gap appear as a mysterious throughput ceiling rather than a clear architectural limitation.

The assistant also assumes that the benchmark methodology is sound—that the random 256-in/256-out prompts are representative, that the concurrency levels are appropriate, and that the measurement tools (bench_serving, decode-batch logs) are accurate. These assumptions are reasonable given the context of previous validation work.

The Broader Implications

This message speaks to a larger pattern in AI infrastructure engineering: the gap between "model support" and "model performance." A model can load, generate correct output, and pass all functional tests while being 10× slower than its potential. This gap is especially pronounced on new hardware architectures where the fast kernel paths are still being developed and gated behind specific compute capabilities.

The assistant's approach offers a template for handling such situations: diagnose thoroughly, document honestly, decide pragmatically, and execute decisively. The diagnosis rules out config issues and identifies the architectural root cause. The documentation preserves the baseline for future comparison. The decision accepts the constraint and refocuses on achievable deliverables. The execution installs NIXL and prepares the PD deployment without delay.

There is also an implicit lesson about the value of custom kernel engineering. The assistant references the earlier K2.6 work where custom CUDA kernels achieved 3-6× speedups. This establishes a precedent: when the off-the-shelf inference stack falls short, targeted kernel development can close the gap. But that work is measured in days, not hours, and must be scoped as a separate project. The message implicitly distinguishes between deployment engineering (configuring and integrating existing components) and kernel engineering (writing new CUDA kernels), and correctly identifies which category the current work falls into.

Conclusion

Message 12381 is a masterclass in honest engineering under constraints. The assistant confronts a 10× performance gap, diagnoses its architectural root cause in the SM120 fallback path, documents the baseline without embellishment, and pivots to the next deliverable without losing momentum. The message captures a moment of technical clarity—the recognition that some bottlenecks cannot be configured away, that hardware architecture imposes hard ceilings, and that the right response is not to fight the ceiling but to work within it while delivering what was promised.

The installation of NIXL that closes the message is a small action with large symbolic weight. It signals that the project is moving forward, that the PD disaggregation architecture will be built and tested, and that the assistant remains focused on the headline deliverable. The SM120 performance story is not finished—it will resurface in later benchmarks, in the PD deployment, and in the final report—but for now, the right call has been made. Measure, document, decide, execute. This is engineering at its most honest and effective.