The Methodical Campaign: Orchestrating Multi-Phase Optimization for DeepSeek-V4-Flash on Blackwell
Introduction
In the high-stakes world of deploying large language models on cutting-edge hardware, the gap between a working deployment and a high-performance one is rarely bridged by a single breakthrough. More often, it is closed through a methodical, multi-phase campaign of measurement, analysis, strategic pivots, and parallel execution—a process that demands not just technical depth but also disciplined workflow management, honest ceiling assessment, and the ability to recognize when the community has already paved the path forward.
This article synthesizes a pivotal segment of an opencode coding session where an AI assistant, working with a user, executed precisely such a campaign. The target: deploying DeepSeek-V4-Flash—a 685B-parameter Mixture-of-Experts model with sparse Multi-head Latent Attention (MLA) and FP4 quantization—on 8× NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture) using the SGLang inference engine. The initial performance was catastrophic: approximately 10 tokens per second at batch size 1 and 25 tok/s at concurrency 16, against a user target of roughly 1000 tok/s—a gap of 40–100×.
What follows is the story of how the assistant navigated from that abysmal baseline through a systematic optimization campaign that spanned upstream integration, kernel autotuning, speculative decoding configuration, and ultimately the recognition that the hardware itself imposed a fundamental ceiling. It is a case study in methodical engineering under uncertainty.
Phase I: Diagnosis and the Hard Ceiling
The campaign began where all serious optimization should: with measurement. A definitive GPU profile captured 40 decode steps totaling 3,830 ms of GPU kernel time and revealed a stark picture. A single kernel—_tiled_sparse_decode_kernel, the sm_120 Triton fallback for sparse MLA attention—consumed 63% of all GPU time (2,412 ms). The MoE GEMV kernel accounted for 10%, FP8 attention projections for 6%, and NCCL all-reduce a mere 2%.
The root cause was architectural. The kernel's grid configuration was (B, H)—at batch size 1, that meant only 64 blocks (1 batch × 64 query heads) launched across approximately 170 streaming multiprocessors (SMs), yielding roughly 38% occupancy. Each block then serially iterated through all 512 top-k tokens in tiles of 16–32. This was the identical low-occupancy pathology that had plagued an earlier Kimi K2.6 deployment, which the assistant had previously fixed with a split-K rewrite achieving 3–6× speedup.
But the assistant's analysis went deeper. Even if the sparse-decode kernel were made instant, the remaining 37% of GPU time—consumed by MXFP4 MoE GEMV, FP8 matmul, and hundreds of small elementwise operations—would limit total speedup to approximately 2.7×, yielding roughly 60 tok/s at concurrency 16. The reason DeepSeek-V4 achieves high throughput on other hardware is the fused DSA/MoE stack—DeepGEMM, trtllm-gen, FP4 C4 indexer—which is arch-gated to SM100 (B200) and simply does not execute on sm_120.
The assistant presented three options to the user: build a bounded split-K sparse-decode kernel (~2–3× overall), undertake a broader owned-sm_120-kernel campaign (multi-week, higher ceiling), or accept the documented ceiling. The user's response was decisive: "Expand the findings with a detailed plan to build the missing kernels tuned for RTX PRO 6000 Blackwell" ([msg 12447]).
Phase II: The Strategic Pivot to Community Intelligence
Before committing to a multi-week kernel engineering effort, the user issued a directive that would fundamentally reshape the campaign: "Also research online in repos if anyone has made issues / PRs about those issues" ([msg 12449]). This single sentence redirected the assistant from first-principles engineering to community-aware strategy.
The research ([msg 12450]) revealed a rich landscape of upstream work that the assistant's local checkout was simply missing:
- SGLang PR #25692 (AliceChenyy): Full SM120 support for DeepSeek-V4/V3 inference, providing Triton-based fallback kernels for all critical paths. The PR explicitly stated that "SM120 desktop Blackwell GPUs lack TMEM, tcgen05, and DeepGEMM support"—directly validating the assistant's diagnosis.
- SGLang PR #20040 (mmangkad): A fix for the MXFP4 MoE
block_ktile size on SM120, resolving a shared-memory budget overflow. - DeepGEMM Issue #317 (rs-ipps): Documentation that
tf32_hc_prenorm_gemmandpaged_mqa_logitskernel families hit "Unsupported architecture" assertions on SM120. - vLLM PR #40991 (jasl): Enables DeepSeek-V4 Flash on SM12x GPUs with a portable Triton sparse MLA path, tested on 2× RTX PRO 6000 Blackwell Workstation Edition. These findings transformed the problem. The assistant was not facing a novel challenge that required bespoke kernel development—it was facing a known gap in its own codebase. The local checkout (commit
735a256from June 17, 2026) contained the base SM120 support from PR #24692 but none of the critical performance follow-ups. The assistant's own git log confirmed this:git log --oneline -40 | grep -iE "sm120|fp8.*config|autotune"returned nothing. The FP8 configs directory contained no RTX PRO 6000-specific files. The assistant's reasoning captured the strategic pivot precisely: "The real insight here is that the path forward isn't necessarily writing new kernels from scratch — it's pulling in the existing upstream performance PRs and configs that are already proven to work" ([msg 12451]). Published benchmarks showed others hitting 37–74 tok/s on similar hardware—4–7× faster than the assistant's current 10 tok/s. The gap was not architectural; it was a gap in applied optimizations.
Phase III: The User's Methodical Directive
The user's response to this analysis was a masterclass in measured decision-making ([msg 12455]):
"update to upstream to pull all those wins, enable MTP, tune kernels (iff upstream tuning can be made better by doing it), set NCCL=LL as that previously was a win (see results in ./), then run and MEASURE so that we can be a bit more methodotical in our next steps attacking the highest impact issues"
This directive contained five distinct actions, each with strategic rationale:
- Update to upstream: Accepting the assistant's analysis that the local checkout was stale, the user authorized pulling the latest SGLang main to capture merged performance improvements.
- Enable MTP: Multi-Token Prediction / EAGLE speculative decoding was the second-highest-ROI quick win. The DeepSeek-V4-Flash model ships a built-in NextN (MTP) layer, making this a configuration change rather than a code change.
- Tune kernels (conditionally): The critical qualifier "iff upstream tuning can be made better by doing it" revealed sophisticated understanding. The user recognized that local autotuning is expensive and may not beat upstream defaults.
- Set NCCL=LL: Grounded in empirical evidence from earlier benchmarks ("see results in ./"), this was a data-driven decision rather than a theoretical one.
- Run and MEASURE: The philosophical core. The user explicitly demanded measurement before further commitment, transforming the conversation from planning to empirical investigation. This directive implicitly answered the assistant's four open questions: yes to Tier A first, maximize sm_120 throughput (not chase the unreachable 1000 tok/s), defer the quantization strategy decision, and keep changes local for now.
Phase IV: Parallel Execution—FP8 Autotuning Meets MTP Research
The assistant translated the user's directive into a concrete execution plan ([msg 12456]). The first action was a targeted diagnostic reconnaissance ([msg 12457])—a single SSH command that checked the git state, searched for FP8 autotune commits, extracted the exact (N,K) shapes needing tuning from server startup logs, confirmed MTP flag availability, and verified the tuning script's presence. The results were revealing: only 4 commits behind upstream (none containing FP8 configs), five distinct FP8 shapes identified, MTP flags confirmed present, and the tuning script at the expected path.
The assistant then pulled the four upstream commits ([msg 12458]), verifying that none touched dependencies or kernel code—the editable install would pick up Python changes automatically. The diff showed 9 files changed, primarily AMD-specific fixes and CI permissions. The pull was a non-event, but the verification was essential risk management.
With the upstream current and the server stopped, the assistant launched the FP8 block-GEMM autotuning campaign ([msg 12461]). The tuning script sweeps through approximately 1,280 candidate configurations per (N,K) shape (combinations of block_m, block_k, block_n, stages, warps, and group sizes) across 18 batch sizes, distributing the work across all 8 GPUs. The assistant created a runner script that loops through the five identified shapes: (1024,4096), (1536,4096), (4096,2048), (4096,512), and (8192,1024). Estimated runtime: 20–40 minutes.
But the assistant did not simply wait. While the tuning ran in the background, it immediately pivoted to researching the MTP configuration ([msg 12462]). The DeepSeek-V4 architecture includes a built-in NextN (multi-token prediction) head, but the exact SGLang launch parameters were unclear. The assistant read the deepseek_v4_hook.py file to understand model-specific constraints, then searched the registered test directory for MTP recipes ([msg 12463]).
The search was methodical. A broad grep across all registered tests for speculative decoding flags returned noisy results from other model families (EAGLE3 for GPT-OSS-120B, Mistral-Large-3). The assistant narrowed to the specific AMD test file test_deepseek_v4_pro_fp4_mtp.py ([msg 12464]), which had been added in one of the pulled upstream commits. A targeted grep revealed the other_args block containing the MTP configuration. The assistant then read the exact lines 78–118 of the test file ([msg 12465]) to extract the complete recipe:
other_args = [
"--trust-remote-code",
"--tp", "8",
"--disable-radix-cache",
"--attention-backend", "dsv4",
# MTP / EAGLE speculative decoding (NextN head from the base model).
"--speculative-algorithm", "EAGLE",
"--speculative-num-steps", "3",
"--speculative-eagle-topk", "1",
"--speculative-num-draft-tokens", "4",
]
Critically, no --speculative-draft-model-path was present—confirming that the NextN head is auto-loaded from the base checkpoint. The assistant adapted this recipe for the Blackwell deployment, choosing the conservative "0xSero recipe" (num-steps=1, num-draft-tokens=2) proven stable on this hardware ([msg 12466]).
Phase V: Orchestrating the Waiting—Benchmark Infrastructure Preparation
The FP8 tuning continued in the background, and the assistant designed its experimental protocol. Two configurations were defined: Config A (baseline optimized: FP8 configs, NCCL LL, CUDA graphs, Marlin MoE) and Config B (identical but with MTP enabled). This clean A/B design isolated MTP as the only variable ([msg 12466]).
The assistant created two scripts ([msg 12468]): serve_dsv4_mtp.sh encoding the MTP launch parameters, and run_measure.sh defining a benchmark harness that tests three concurrency levels (C=1, 8, 16) with 256-token sequences, capturing throughput, TPOT, TTFT, and acceptance length. The decision to drop --num-continuous-decode-steps from both configurations was deliberate—earlier measurements showed it had no effect, and including it would only add noise.
The tuning progress was monitored through a polling loop ([msg 12469]) that checked two independent signals every 90 seconds: the count of PRO_6000 config files written to the configs directory, and the presence of "TUNE_DONE" in the tuning log. The output over 15 minutes showed the expected pattern: the first shape (N=1024, K=4096) completed in ~6 minutes, followed by the second shape (N=1536, K=4096) in another ~9 minutes. The total was projected at 30–50 minutes, within the assistant's estimate.
The Broader Significance: Methodical Optimization as Engineering Discipline
This chunk of the conversation is remarkable not for a single dramatic breakthrough but for the sustained methodological discipline it displays across dozens of messages. Several patterns emerge that are worth studying:
Measurement before action: Every significant decision was preceded by measurement. The GPU profile identified the bottleneck. The git log check revealed the missing upstream commits. The server log grep extracted the exact FP8 shapes. The test file read confirmed the MTP recipe. The assistant never guessed—it always checked.
Strategic pivots grounded in evidence: The assistant's initial plan was to build custom kernels. The user's request to research online revealed that upstream integration was a faster path. The assistant pivoted without ego, recognizing that community work had already solved parts of the problem. This is the hallmark of mature engineering: the willingness to change direction when evidence demands it.
Parallel workflow orchestration: The assistant consistently overlapped long-running compute tasks with preparation for subsequent phases. While FP8 tuning ran, it researched MTP configuration. While waiting for tuning to complete, it wrote benchmark scripts. This pipeline approach minimized the critical path and maximized productivity during inevitable waiting periods.
Honest ceiling communication: Throughout the campaign, the assistant maintained intellectual honesty about what was achievable. The 1000 tok/s target was not reachable on sm_120 hardware—that required B200 or SM100-class GPUs. The assistant stated this clearly and repeatedly, preventing wasted effort chasing an impossible target.
Community-aware engineering: The single most impactful decision in this chunk was the user's instruction to research existing PRs and issues. This transformed the campaign from a solo engineering effort into an integration project, leveraging work from AliceChenyy, mmangkad, rs-ipps, jasl, and others across SGLang, vLLM, and DeepGEMM.
Conclusion
The optimization campaign captured in this chunk is a case study in how to systematically improve LLM inference performance on novel hardware. The assistant moved from a 10 tok/s baseline through diagnosis, strategic pivoting, upstream integration, kernel autotuning, and speculative decoding configuration—all while maintaining methodological discipline and honest communication about what the hardware could deliver.
The FP8 autotuning would eventually yield a modest ~6% improvement (since FP8 GEMM is only a small fraction of decode time). MTP would provide a 47% boost at single-request throughput but zero gain at concurrency due to verifier saturation. The NVFP4 quantization switch would route MoE through tensor cores for a ~24% improvement. But each of these gains was measured cleanly, attributed correctly, and built upon a foundation of methodical engineering.
The user's directive to "run and MEASURE" before committing to the next phase proved prescient. It ensured that every optimization lever was evaluated on its actual impact, not its theoretical potential. And when the measurements revealed that configuration tuning could only deliver incremental gains against a structural sm_120 bottleneck, the assistant had the data to make that case convincingly.
In the end, this chunk is a testament to the power of methodical, measurement-driven optimization—and a reminder that in systems engineering, the most valuable skill is knowing when to measure, when to pivot, and when to tell the hard truth about what the hardware can deliver.