The 24% Ceiling: When Configuration Tuning Meets the sm_120 Kernel Wall

Introduction

In the high-stakes world of large language model deployment, the gap between a model that runs and one that performs can span orders of magnitude. 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.

What follows is a methodical, measurement-driven optimization campaign that tested every configuration lever available within the SGLang framework — NCCL communication tuning, CUDA graph capture, MTP speculative decoding, FP8 GEMM autotuning, tilelang indexer fusion, expert parallelism, and a pivotal switch from stock MXFP4 to NVIDIA's NVFP4 quantization format. Each optimization delivered measurable improvement — 6% here, 24% there, 47% for single requests — but the aggregate still fell far short of the hardware's roofline potential. The decisive bottleneck, revealed through GPU profiling and confirmed through cross-validation, was the sm_120 fallback kernel path: the sparse-MLA attention and MoE slot-GEMV kernels that should run on tensor cores were instead executing on CUDA cores via Triton fallbacks, consuming 63% and 10% of decode time respectively. The NVFP4 checkpoint fixed the MoE path, routing it through tensor cores, but attention remained the primary limiter — and the MTP verifier's memory consumption silently halved the effective batch size.

This is a case study in the limits of configuration optimization, the importance of measurement-driven iteration, and the moment when an engineering team must recognize that the remaining gap requires custom kernel development rather than flag tuning.

The Pivot: From Custom DDTree Kernels to DeepSeek-V4-Flash

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 that achieved 3–6× decode speedup over Triton, and deploying CUDA graph capture-safety for real-serving wins. 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 ([msg 12348]).

The assistant's response was methodical. Rather than diving into implementation, it launched a parallel reconnaissance phase, 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." 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.

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. The editable install of SGLang required careful dependency management — resolving import shadowing from leftover site-packages directories and debugging a shebang path issue in the launch script. After these hurdles, the model loaded successfully on a single-node TP4 configuration and generated correct output: "The capital of France is Paris."

Prefill-Decode Disaggregation: The Architectural Deliverable

The headline deliverable — prefill-decode disaggregation — was then orchestrated ([msg 12385]). 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.

The architecture was sound. The model loaded correctly — 146 GB of FP4-quantized weights across 46 shards. The servers reported "fired up and ready to roll" within 60 seconds, benefiting from a warm JIT cache. An end-to-end test produced a coherent answer about prefill-decode disaggregation itself — a self-referential validation that the entire pipeline functioned correctly.

But the benchmark results, when they arrived after several minutes of polling, were sobering:

| Configuration | Throughput (tok/s) | TPOT (ms) | TTFT (ms) | |---|---|---|---| | Single-node TP4, bs=1 | 10.2 | 94 | 696 | | PD disaggregation, C=1 | 9.95 | 94 | 893 | | PD disaggregation, C=8 | 24.5 | 282 | 2,000 | | PD disaggregation, C=16 | 25.0 | 510 | 8,800 |

The pattern was unmistakable: throughput plateaued at ~25 tok/s regardless of concurrency, and time-to-first-token degraded catastrophically as queue buildup overwhelmed the decode side. The assistant's diagnosis was precise: "PD doesn't raise aggregate tok/s here — decode is the bottleneck, not prefill interference or KV transfer." Prefill-decode disaggregation, for all its architectural elegance, cannot help when the decode phase itself is the binding constraint.

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 User's Challenge and the Diagnostic Pivot

The assistant's initial conclusion — that the sm_120 fallback kernels represented a hard hardware ceiling — was met with a firm rejection from the user. In [msg 12395], the user wrote: "We expect much much faster than 25T/s on this model, at C=16 should be at/above 1k tps." This was not a suggestion; it was a statement of fact about what the hardware should be capable of. The user pointed to two specific levers: NCCL LL (low-latency protocol) and alternatives to the Marlin MoE backend, hinting that optimizations previously applied in the repository (from work on the Kimi K2.6 model) were missing from the current deployment.

This intervention fundamentally reshaped the campaign ([msg 12396]). The assistant pivoted from acceptance to investigation, launching parallel searches of the local repository for prior optimization work and querying the remote server's environment configuration. The research was decisive: the K2.6 service had achieved 1,291 tok/s with TP8 on the same hardware using a specific NCCL PCIe tuning block — and the current DeepSeek-V4-Flash deployment had none of those settings. The NCCL environment variables that should have been set — NCCL_PROTO=LL, NCCL_ALGO=Ring, NCCL_P2P_LEVEL=5, NCCL_MIN_NCHANNELS=8, NCCL_MAX_NCHANNELS=16, NCCL_BUFFSIZE=16777216, NCCL_NTHREADS=512 — were all absent.

The user reinforced this direction with a four-word message in [msg 12397]: "cuda graph too, especially." This terse correction elevated CUDA graphs — which had delivered a documented 3.8× speedup on K2.6 (from 26 to 98 tok/s) — above other optimizations in priority. The assistant recognized that "CUDA graphs enabled" and "CUDA graphs working effectively" are not the same thing, and that the current deployment might be falling back to eager execution for uncaptured batch sizes.

Message [msg 12398] captures the synthesis of this research into a concrete action plan. The assistant created a shared NCCL environment file (dsv4_nccl_env.sh) encoding the proven PCIe tuning parameters, and an updated server launch script (serve_dsv4_tp4.sh) with --cuda-graph-max-bs 32, --num-continuous-decode-steps 4, and --mem-fraction-static 0.85. The strategy was to validate on single-node TP4 first — the simplest configuration — before applying the optimizations to the disaggregated deployment.

But the execution hit a snag. The pkill command intended to stop the old PD servers was followed by a nohup relaunch, but the new server never started. The log file showed stale content from the previously killed process, and the assistant fell into a monitoring loop, checking every 30 seconds for a server that would never arrive ([msg 12399]). The user's single-word intervention — "proceed" ([msg 12400]) — broke the loop, redirecting effort from observation to diagnosis.

The subsequent investigation ([msg 12401]) revealed the root cause: the old prefill process (PID 73216) had survived the pkill command, still holding port 30000. The new server, unable to bind to the port, had failed silently — a classic race condition where asynchronous process termination and background shell semantics conspired to produce a silent failure. The assistant's five-dimensional state check — processes, scripts, logs, GPU memory, and health endpoint — provided a reusable debugging methodology for future deployment attempts.

The debugging continued through several more rounds: a failed SSH command due to quoting issues ([msg 12402]), a server that started but was immediately killed by its own pkill command matching itself ([msg 12403], [msg 12404]), and an OOM crash caused by insufficient memory for the model weights plus NCCL buffers ([msg 12407], [msg 12410]). Each failure was diagnosed and addressed: the self-kill was fixed by using kill with a specific PID instead of pkill -f; the OOM was resolved by reducing --mem-fraction-static from 0.85 to 0.70 and ensuring the NCCL environment was properly sourced.

The Optimization Campaign: Incremental Gains Against a Ceiling

Once the server was stable, the assistant executed a systematic, measurement-driven optimization campaign. Each lever was benchmarked individually, with results documented and analyzed before proceeding to the next.

FP8 GEMM autotuning. The assistant pulled the latest upstream SGLang and generated SM120 FP8 GEMM autotune configurations. The result: a ~6% improvement in throughput. This was consistent with profiling data showing that FP8 GEMM accounts for only ~6% of decode time — optimizing a small fraction of the workload yields proportionally small gains.

NCCL low-latency tuning. The NCCL LL protocol and Ring algorithm, which had been transformative for K2.6, produced negligible improvement for DeepSeek-V4-Flash. Profiling revealed why: communication accounts for only ~2% of decode time on this configuration. The NCCL tuning, while essential on systems where all-reduce dominates, was attacking a non-bottleneck.

MTP/EAGLE speculative decoding. This was the most dramatic single-request optimization: MTP boosted throughput by 47% for individual requests, from ~10 tok/s to ~15 tok/s. However, at concurrency, the gain vanished. The MTP verifier — the draft model that validates speculative tokens — became saturated, consuming enough GPU memory to effectively halve the batch size from 16 to 8. The verifier's memory footprint meant that at C=16, only 8 slots were actually doing productive decode work, while the other 8 were consumed by verifier state.

Tilelang indexer fusion. The assistant attempted to use tilelang's JIT compilation to fuse the sparse attention indexer operations ([msg 12425]). The tilelang indexer path (SGLANG_OPT_USE_TILELANG_INDEXER) is checked before the forced torch path in the decision tree, and tilelang 0.1.8 was already installed. However, when the server was restarted with the flag enabled, it immediately crashed — the tilelang compiler failed during JIT compilation on sm_120 with CUDA 13 ([msg 12430]). The error was a compile-time failure, not a runtime issue, confirming that tilelang's kernel generation pipeline does not support the sm_120 architecture.

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. Switching to a non-Marlin backend would produce incorrect results by using the wrong quantization format.

Expert Parallelism (EP). The assistant attempted to split experts across GPUs to reduce the per-GPU MoE computation ([msg 12434], [msg 12435]). The EP4 configuration launched successfully but performed worse than TP4: C=1 throughput dropped to 8 tok/s (vs 10), and C=16 dropped to 14 tok/s (vs 23). The PCIe all-to-all overhead for expert routing dominated the compute time at these small batch sizes, making EP counterproductive.

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. CUDA graphs remove Python launch overhead but cannot accelerate kernel execution time itself.

The NVFP4 Pivot: Routing MoE Through Tensor Cores

The most impactful change was switching from the stock MXFP4 checkpoint to NVIDIA's official NVFP4 quantization format ([msg 12418]). This was not a trivial change — it required fetching PR #25820 for correct NVFP4 auto-detection and downloading a 149 GB checkpoint. The NVFP4 format routes MoE execution through tensor-core paths (Marlin W4A16 or native cutlass FP4 grouped GEMM), whereas the W4A16 path fell through to Triton fallback.

Both NVFP4 backends delivered identical throughput of ~28 tok/s at C=16, a 24% improvement over the baseline. Crucially, this confirmed that the MoE path was now executing on tensor cores rather than CUDA cores. The NVFP4 checkpoint had fixed the MoE path — but attention remained the primary limiter.

The assistant also discovered a critical hidden factor: the MTP verifier's memory consumption was silently halving the effective batch size. At C=16, the verifier consumed enough GPU memory that only 8 requests could be processed simultaneously. This explained why throughput plateaued at ~25 tok/s regardless of concurrency — the system was effectively running at C=8, not C=16.

The user's roofline analysis confirmed that the hardware should be capable of 300–600 tok/s at C=16. With 1.9 TB/s of VRAM bandwidth and 13B active parameters, the theoretical peak was orders of magnitude above the measured throughput. The gap was not in the hardware — it was in the software path.

The Decisive Profile: 63% in One Kernel

The assistant ran a definitive GPU profile to understand where the time was going ([msg 12441], [msg 12442]). The trace captured 40 decode steps with the torch profiler, producing compressed JSON trace files for all four TP ranks. Parsing the TP-0 trace revealed the definitive bottleneck:

| Kernel | Total Time | Share | Calls | |---|---|---|---| | _tiled_sparse_decode_kernel | 2412 ms | 63% | 3276 | | _mxfp4_slot_gemv_kernel (MoE) | 380 ms | 10% | 3354 | | _w8a8_block_fp8_matmul (attn proj) | 244 ms | 6% | 9204 | | ncclAllReduce (comm) | 85 ms | 2% | 3393 |

The decisive finding was that 63% of decode time was consumed by 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 in tiles of 16-32. The result is catastrophic under-utilization: the majority of GPU compute units sit idle while a tiny fraction do the work ([msg 12444]).

The assistant confirmed the root cause by examining the kernel source ([msg 12444]). Line 246 of flash_mla_sm120_triton.py reads: grid = (B, H). At batch size 1, that's 64 blocks total on ~170 SMs — approximately 38% occupancy, with a long serial inner loop over 512 top-k tokens. There is no top-k/KV-split dimension in the grid. This is the classic low-occupancy bs=1 pathology — identical to the K2.6 verify kernel before split-K parallelization was applied.

The assistant documented these findings in a comprehensive PROFILE_FINDINGS.md file ([msg 12445]), capturing the precise bottleneck, the ceiling analysis, and the path forward. The ceiling analysis was sobering: even a perfect (0ms) sparse-decode kernel would only reduce total GPU time from 3830ms to 1418ms — a 2.7× improvement, yielding ~60 tok/s at C=16. The remaining 37% (MXFP4 MoE GEMV, FP8 matmul, elementwise ops) are also sm_120 fallbacks. The fast fused DSA/MoE stack (DeepGEMM, trtllm-gen, FP4 indexer) is arch-gated to SM100 and simply does not run on sm_120.

The Path Forward: Custom Kernel Engineering

The campaign reached a natural inflection point. Configuration tuning — NCCL protocols, CUDA graphs, memory fractions, continuous decode steps — had been exhausted. Checkpoint format changes (MXFP4 to NVFP4) had delivered their ~24% gain. MTP speculative decoding had shown promise for single requests but was neutralized at concurrency by verifier memory overhead.

The remaining bottleneck was structural: the sm_120 fallback kernels for sparse-MLA attention. The fast fused paths (DeepGEMM, trtllm-gen, FP4 indexer) are architecture-gated to SM100 and do not run on sm_120. The only path forward is custom kernel development — the same playbook that delivered 3–6× gains for K2.6 through a custom verify-attention kernel.

Two concrete paths emerged:

1. Disable MTP to restore the full batch capacity of 16, then re-measure to establish a clean baseline without verifier overhead. 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 47% single-request speculative decoding benefit, resulting in roughly equivalent performance at low concurrency.
2. Build a split-K tensor-core sparse-attention kernel for sm_120, attacking the 63% attention bottleneck directly. The fix is a split-K strategy: instead of grid (1, 64), reshape to (1, 64, 8) = 512 blocks, which would give full occupancy and potentially 3–6× speedup on this kernel alone. Since the sparse decode kernel accounts for 63% of total time, that could translate to a ~2× overall improvement. This is exactly the approach that proved transformative for K2.6. 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.

The hidden cost of speculative decoding. The MTP verifier's memory consumption silently halved the effective batch size, neutralizing the throughput gains at concurrency. This is a cautionary tale about the hidden costs of speculative decoding: the verifier is not free, and its memory footprint can negate the benefits of increased batch capacity.

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 Blackwell Ceiling: Deploying DeepSeek-V4-Flash with Prefill-Decode Disaggregation and Confronting the sm_120 Fallback Kernel Bottleneck" — Article on the initial deployment and PD disaggregation architecture.

[2] "The Methodical Pursuit of Performance: Optimizing DeepSeek-V4-Flash on Blackwell GPUs" — Article on the systematic optimization campaign.

[3] "The Methodical Campaign: Orchestrating Multi-Phase Optimization for DeepSeek-V4-Flash on Blackwell" — Article on the orchestration of the optimization phases.

[4] "The 24% Ceiling: When Configuration Tuning Meets the sm_120 Kernel Wall" — Article on the NVFP4 quantization pivot and the structural bottleneck diagnosis.