The Turning Point: From Custom Kernel Rewrites to Upstream Integration

Introduction

In the high-stakes world of deploying large language models on cutting-edge hardware, few moments are as decisive as the one captured in message 12451 of this opencode session. The assistant, having spent hours methodically profiling, tuning, and diagnosing a DeepSeek-V4-Flash deployment on 8× NVIDIA RTX PRO 6000 Blackwell GPUs, suddenly confronts a revelation that reshapes the entire optimization strategy. What begins as a routine check of upstream repositories transforms into a fundamental reframing of the problem: the path forward is not about writing custom CUDA kernels from scratch, but about integrating existing, proven optimizations that the assistant's checkout was simply missing.

This message represents a critical inflection point in a multi-day engineering effort. The session had been chasing a throughput target of 1000 tokens per second—a figure that seemed increasingly unattainable as every configuration lever was pulled and every tuning knob turned, only to yield marginal gains. The assistant had already identified the primary bottleneck: a Triton-based sparse-MLA decode kernel consuming 63% of GPU time, running at only ~38% occupancy on the 188-streaming-multiprocessor Blackwell GPU. The proposed solution was a split-K kernel rewrite, following the playbook that had delivered 3-6× speedups for a similar bottleneck in an earlier K2.6 deployment. But the user's request to "expand the findings with a detailed plan to build the missing kernels" prompted a crucial detour: researching what the open-source community had already discovered about running DeepSeek-V4 on SM120 hardware.

What the assistant found would change everything.

The Message in Full

The message opens with an extended reasoning section that reveals the assistant's thinking process in real time:

Looking at the upstream research, there's a rich landscape of SGLang PRs and issues around SM120 support for DeepSeek-V4, including the fallback kernels we're currently using and some promising FP8 GEMM autotune configurations that are showing significant throughput improvements.

>

But I'm noticing something critical: the profiling breakdown from the PR #24692 author shows a very different picture than what I'm seeing. They report MoE at 40% of time-per-output-token, communication overhead at 20-25%, dense FP8 GEMM at 15-20%, and sparse decode only 10-12ms — yet my profile shows sparse decode consuming 63% of the time. That's a major discrepancy I need to reconcile, especially since we're on different tensor parallelism configurations.

>

The difference could stem from several factors: my TP4 setup versus their TP8 and TP2 runs, the fact that my profiling window captured both prefill and decode phases with multiple sparse attention launches per layer, or possibly over-attribution in how the kernel calls were counted. Regardless of which component is actually the primary bottleneck, both analyses point to the same conclusion — the slowdown is distributed across sparse MLA decode, MXFP4 MoE, FP8 GEMM, and NCCL communication. The FP8 GEMM autotune from PR #25696 is still a proven quick win delivering 4-5x speedup, pushing throughput to 44-52 tokens/second on TP2.

>

Comparing against published numbers reveals I'm running 4-7x slower than the SM120 baseline — others are hitting 40-75 tokens/second on dual RTX PRO 6000 while I'm stuck at 10 tokens/second on TP4. This gap strongly suggests I'm missing the upstream optimizations, particularly the FP8 GEMM autotune configs from PR #25696, which my current branch doesn't have (the warning confirms I'm using the default W8A8 kernel config instead of the SM120-specific one). Looking at the progression of published results, my 10 tokens/second at batch size 1 aligns with the baseline from PR #24047 before the major optimizations were applied, indicating I have the foundation but lack the critical tuning layers that others have already integrated.

>

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. Additionally, both the 0xSero and jasl results leverage EAGLE/MTP speculative decoding for 1.5-2x speedup, which I'm not using despite the model having the MTP layer available. Even with all these optimizations applied, the realistic ceiling for aggregate throughput at C=16 on this hardware is probably 40-100 tokens/second based on what the community has published, so the user's 1000 token/second expectation isn't achievable — but getting from 10 to 40-50+ tokens/second through existing optimizations is a solid 4-5x win without custom kernel work.

>

The concrete next steps are checking whether the FP8 autotune configs from the upstream PR are already in my checkout, enabling MTP/EAGLE speculative decoding with the proven recipe, and potentially pulling in the DeepGEMM SM120 kernels that add the fused operations. There's also some conflicting information about NCCL settings — the published results use different configurations (Tree vs Ring, P2P disabled vs enabled) so I'll need to test which actually works on this hardware.

The reasoning continues, exploring the discrepancy in profiling data and the implications for the optimization strategy:

I'm noticing a key discrepancy in the profiling data: my profile shows sparse decode at 63%, but upstream reports MoE at 40%. The difference likely comes down to batch size — at bs=1 (my profiling), sparse decode dominates because attention is the bottleneck, but at higher concurrency like bs=8+, MoE becomes the limiting factor. Both are true depending on the workload, so the plan needs to address both. For the C=16 target aiming for 1000 tok/s, MoE optimization matters more since we'll be running at higher batch. I should re-profile per-component after Tier A to confirm where the actual bottlenecks are at production concurrency.

>

Now I'm checking what's actually achievable: published benchmarks show nobody exceeds 50-100 tok/s aggregate on 2-4x RTX PRO 6000 for this model. Getting to 1000 tok/s at C=16 would require 62.5 tok/s per request, but at that concurrency level the MoE saturates and per-request throughput drops significantly. Even with a proper grouped tensor-core MoE implementation, I'm realistically looking at maybe 100-200 tok/s aggregate at C=16 with full Tier B optimizations. The honest assessment is that 1000 tok/s on this hardware for this model isn't feasible — that would require a B200 or a smaller/denser model. I need to set clear expectations with the user about what's actually achievable.

After the reasoning, the assistant summarizes the upstream discoveries and executes two verification commands:

Rich upstream landscape found — this substantially sharpens the plan. Key discoveries: SGLang PR #25696 (SM120 FP8 GEMM autotune → 44-52 tok/s TP2, 4-5×), the jasl/DeepGEMM sm120 branch + DeepGEMM PR #318/#319 (the missing mHC + paged-MQA-logits kernels), vLLM jasl PRs #40991/#41834 (split-KV sparse MLA + MTP), and lna-lab/QUTLASS SM120 NVFP4 GEMM references. Also: published SM120 numbers are 37-74 tok/s single-req with MTP — I'm getting 10 because I'm missing those layers. Let me verify what's already in my checkout and the exact device, then present the plan.

The first bash command checks the local git repository for relevant commits and FP8 configuration files:

cd /home/theuser/glm-kimi-sm120-rtx6000bw/sglang-dsv4 && echo "=== HEAD date ===" && git log -1 --format="%h %ci" && echo "=== recent sm120/dsv4/fp8-config commits in checkout ===" && git log --oneline -40 | grep -iE "sm120|sm_120|deepseek.v4|dsv4|fp8.*config|autotune|mxfp4|mqa|sparse|tilelang" | head -25 && echo "=== any RTX PRO 6000 fp8 configs present? ===" && ls python/sglang/srt/layers/quantization/configs/ | grep -iE "PRO_6000|Blackwell|RTX" || echo "NONE (need to generate)"

The output confirms the suspicion:

=== HEAD date ===
735a256 2026-06-17 21:18:41 +0800
=== recent sm120/dsv4/fp8-config commits in checkout ===
=== any RTX PRO 6000 fp8 configs present? ===
NONE (need to generate)

The second command queries the remote GPU server for hardware specifications:

timeout 30 ssh -o StrictHostKeyChecking=no root@10.1.230.171 '/root/venv_sglang211/bin/python -c "import torch; p=torch.cuda.get_device_properties(0); print(\"name:\",p.name); print(\"SMs:\",p.multi_processor_count); print(\"smem/block optin KB:\", p.shared_memory_per_block_optin/1024 if hasattr(p,\"shared_memory_per_block_optin\") else \"n/a\"); print(\"cc:\", p.multi_processor_count); print(\"mem GB:\", round(p.total_memory/1e9,1))"'

The output confirms the hardware:

name: NVIDIA RTX PRO 6000 Blackwell Server Edition
SMs: 188
smem/block optin KB: 99.0
cc: 12 0
mem GB: 102.0

Why This Message Was Written: Reasoning, Motivation, and Context

To understand why this message exists, we must trace the arc of the session that produced it. The assistant had been engaged in a multi-day effort to deploy DeepSeek-V4-Flash on a machine equipped with 8× NVIDIA RTX PRO 6000 Blackwell GPUs. This is a formidable hardware configuration: each GPU offers 102 GB of VRAM, 188 streaming multiprocessors, and support for the SM120 compute capability (compute capability 12.0). The Blackwell architecture introduces new tensor core features (SM120a with FP4/FP8 tensor cores) but crucially lacks the TMEM, tcgen05, and DeepGEMM support that make the B200 (SM100) so performant for this model.

The user's target was ambitious: 1000 tokens per second at concurrency 16. The assistant had methodically worked through every available configuration lever—NCCL protocol tuning, CUDA graphs, expert parallelism, tilelang indexer fusion, non-Marlin MoE backends—and none had moved throughput beyond ~23 tok/s at C=16. A definitive GPU profile had traced 63% of decode time to a single kernel: _tiled_sparse_decode_kernel, the SM120 Triton fallback for sparse MLA attention, which launched only 64 blocks on ~170 SMs, serially iterating all 512 top-k tokens.

The assistant had documented this bottleneck and proposed a split-K kernel rewrite, following the successful playbook from an earlier K2.6 deployment where a similar fix had delivered 3-6× speedup. But the user, rather than approving the rewrite, asked the assistant to "expand the findings with a detailed plan to build the missing kernels tuned for RTX PRO 6000 Blackwell" (msg 12449) and to "research online in repos if anyone has made issues / PRs about those issues" (msg 12449).

This message is the assistant's response to that request. It represents the convergence of two investigative threads: the deep technical understanding gained from hours of profiling and the new external knowledge from upstream research. The motivation is to produce a comprehensive, evidence-based plan that accounts for what the community has already discovered, avoiding the trap of reinventing wheels that others have already built.

How Decisions Were Made

The decision-making process in this message is particularly instructive because it involves a fundamental strategic pivot. The assistant had been operating under the assumption that the path forward required custom kernel development—specifically, a split-K rewrite of the sparse-MLA decode kernel. This assumption was reasonable given the profiling data, which showed that kernel consuming 63% of GPU time. But the upstream research revealed a different picture.

The key decision was to reconcile the discrepancy between local profiling and published results. The assistant noticed that PR #24692 reported a very different breakdown: MoE at 40%, communication at 20-25%, dense FP8 GEMM at 15-20%, and sparse decode at only 10-12ms. This discrepancy could not be ignored. The assistant considered several explanations: different tensor parallelism configurations (TP4 vs TP8/TP2), profiling window artifacts (capturing both prefill and decode), or over-attribution in kernel counting. But the crucial insight was that both analyses pointed to the same conclusion: the slowdown was distributed across multiple components, and the FP8 GEMM autotune from PR #25696 was a proven quick win delivering 4-5× speedup.

This led to a second decision: prioritize upstream integration over custom kernel development. The assistant realized that the 4-7× performance gap to published SM120 baselines was not due to fundamental architectural limitations but to missing optimization layers—specifically the FP8 GEMM autotune configs that PR #25696 provided. The local checkout confirmed this: git log showed no relevant sm120/dsv4/fp8-config commits, and the FP8 configs directory contained no RTX PRO 6000-specific files.

A third decision was to set realistic expectations about the throughput ceiling. The assistant calculated that even with all optimizations applied, the realistic aggregate throughput at C=16 on this hardware was 40-100 tok/s, not the 1000 tok/s target. This was a crucial reality check, grounded in published benchmarks showing nobody exceeding 50-100 tok/s on 2-4× RTX PRO 6000 GPUs for this model. The assistant recognized that the 1000 tok/s target would require a B200 or a smaller/denser model—a hard architectural constraint, not a tuning problem.

Assumptions Made by the Assistant

Several assumptions underpin the reasoning in this message, some explicit and some implicit:

1. The profiling discrepancy is due to batch size differences. The assistant hypothesizes that at bs=1, sparse decode dominates (63%), but at higher concurrency (bs=8+), MoE becomes the bottleneck. This is a reasonable hypothesis but unverified—the assistant explicitly notes the need to re-profile after Tier A optimizations to confirm.

2. The upstream PRs are compatible with the current codebase. The assistant assumes that pulling in PR #25696 (FP8 GEMM autotune), the jasl/DeepGEMM sm120 branch, and the vLLM PRs will integrate cleanly. This is an optimistic assumption; in practice, dependency conflicts, API changes, and version mismatches often complicate such integrations.

3. MTP/EAGLE speculative decoding will provide 1.5-2× speedup. This assumption is based on published results from 0xSero and jasl, but the assistant hasn't verified that the model's MTP layer is compatible with the current SGLang version or that it works correctly on SM120 hardware.

4. The hardware ceiling is 40-100 tok/s at C=16. This is extrapolated from published benchmarks on 2-4× RTX PRO 6000 GPUs, but the target machine has 8 GPUs. The assistant doesn't fully explore whether the additional GPUs could enable higher throughput through better parallelism strategies (e.g., TP8 instead of TP4).

5. The FP8 GEMM autotune configs are the primary missing optimization. The assistant's git log search shows no relevant commits, but this doesn't prove the configs are absent—they could be generated at runtime or loaded from a different path. The assumption is validated by the directory listing showing no RTX PRO 6000-specific files, but this is still an inference.

Mistakes or Incorrect Assumptions

While the message is analytically sound, several elements warrant critical examination:

1. The profiling discrepancy may be more fundamental than batch size. The assistant attributes the 63% vs 10-12ms sparse decode discrepancy to batch size differences, but the upstream PR #24692 reports "sparse decode only 10-12ms" in absolute terms, not as a percentage. If the upstream total time-per-token is much lower (e.g., 30-50ms total), then 10-12ms for sparse decode would be 20-40%—still significant but not dominant. The assistant's own profile shows 2412ms for sparse decode out of 3830ms total (63%). The difference in absolute terms (2412ms vs 10-12ms) is so vast that batch size alone cannot explain it. This suggests a more fundamental issue: perhaps the assistant's profiling methodology overcounts kernel invocations (e.g., counting multiple sparse attention calls per layer as separate invocations when they should be amortized), or the upstream numbers are for a different model configuration (e.g., shorter context, fewer top-k tokens).

2. The assumption that upstream integration is easier than custom kernels may be optimistic. The assistant proposes pulling in PRs from SGLang, vLLM, and DeepGEMM—three separate projects with different APIs, versioning, and dependency trees. The SGLang checkout is at a specific commit (735a256 from June 17, 2026), and the upstream PRs may depend on later changes or have conflicts. The assistant doesn't assess the integration complexity or risk of regressions.

3. The hardware ceiling estimate may be too conservative. The assistant estimates 40-100 tok/s at C=16 based on published 2-4× GPU results, but the target machine has 8 GPUs. If the optimizations scale with GPU count (which is not guaranteed—communication overhead grows with TP degree), the ceiling could be higher. Conversely, the assistant's earlier EP4 experiments showed worse performance due to PCIe all-to-all overhead, suggesting that 8-GPU scaling may hit diminishing returns.

4. The message doesn't fully address the user's request for a "detailed plan to build the missing kernels." The user explicitly asked for a kernel-building plan, but the assistant's response pivots toward upstream integration. While this pivot is justified by the research findings, the assistant doesn't provide the kernel-building plan the user requested—it defers that to a future response ("Let me verify what's already in my checkout and the exact device, then present the plan"). The message ends with verification commands but no plan document.

Input Knowledge Required to Understand This Message

To fully grasp this message, the reader needs substantial domain knowledge spanning multiple areas:

Hardware Architecture: Understanding the NVIDIA Blackwell GPU architecture, specifically the distinction between SM120 (desktop/workstation Blackwell: RTX PRO 6000, RTX 5090) and SM100 (datacenter Blackwell: B200). The critical detail is that SM120 lacks TMEM, tcgen05, and DeepGEMM support, forcing reliance on Triton fallback kernels for operations that B200 handles natively through fused tensor-core paths.

Model Architecture: DeepSeek-V4-Flash is a Mixture-of-Experts (MoE) model with Multi-head Latent Attention (MLA) and sparse attention mechanisms. The "sparse decode" kernel handles the MLA attention computation with top-k token selection. The model also uses MXFP4 quantization for MoE weights and FP8 for attention projections. Understanding the interaction between these components is essential for interpreting the profiling data.

Inference Framework Architecture: The message references SGLang's kernel dispatch mechanism, specifically how it selects between Triton fallback kernels (for SM120) and native CUDA kernels (for SM100). The FP8 GEMM autotune configs are JSON files that specify tile sizes, pipeline stages, and other parameters for the W8A8 block-scaled matrix multiplication kernel.

Profiling and Performance Analysis: The message relies on NVIDIA GPU profiling to attribute kernel time. Understanding the difference between kernel-level profiling (what the assistant did) and time-per-output-token reporting (what the upstream PRs did) is crucial for reconciling the discrepancy.

Version Control and Dependency Management: The assistant's git log check and directory listing reveal the state of the local checkout. The reader must understand that the absence of certain commits and config files implies missing optimizations, and that pulling upstream PRs requires careful integration.

Speculative Decoding: MTP (Multi-Token Prediction) and EAGLE are speculative decoding techniques that predict multiple tokens per forward pass, trading compute for latency. The assistant assumes these will provide 1.5-2× speedup, which requires understanding how they interact with the model architecture and hardware.

Output Knowledge Created by This Message

This message creates several forms of valuable knowledge:

1. A verified gap analysis between local deployment and published baselines. The assistant confirms that the local checkout is missing FP8 GEMM autotune configs (no RTX PRO 6000-specific files in the configs directory) and relevant optimization commits (empty git log grep). This transforms a suspicion into a confirmed diagnosis.

2. A catalog of upstream resources for SM120 DeepSeek-V4 optimization. The message identifies five key resources:

The Thinking Process: A Window into Engineering Decision-Making

The extended reasoning section of this message is a rare and valuable artifact: a real-time record of an experienced engineer's thought process as they synthesize new information and adjust their strategy. Let me analyze its structure and methodology.

Phase 1: Pattern Recognition. The assistant begins by noting the "rich landscape" of upstream work, immediately recognizing that the community has already invested significant effort in SM120 DeepSeek-V4 optimization. This is not a blank-slate problem—there are existing solutions to leverage.

Phase 2: Discrepancy Detection. The assistant notices the profiling discrepancy between local results (63% sparse decode) and upstream reports (40% MoE). This is the critical moment—the assistant could have dismissed the upstream data as incomparable (different TP degree, different profiling methodology), but instead treats it as a signal that demands investigation.

Phase 3: Hypothesis Generation. The assistant generates multiple hypotheses for the discrepancy: TP4 vs TP8/TP2, profiling window artifacts, over-attribution. This is a mature engineering response—rather than seizing on a single explanation, the assistant considers multiple possibilities and notes that "both analyses point to the same conclusion."

Phase 4: Quantitative Comparison. The assistant compares local throughput (10 tok/s) against published baselines (37-74 tok/s) and calculates the gap (4-7×). This quantification transforms a vague feeling of "slow performance" into a precise, actionable metric.

Phase 5: Root Cause Inference. The assistant connects the performance gap to the missing FP8 GEMM autotune configs, noting that the local throughput aligns with the pre-optimization baseline from PR #24047. This is a powerful inference: the assistant doesn't just know that performance is bad—it knows why it's bad and what specific optimization is missing.

Phase 6: Strategic Reframing. The assistant articulates the key insight: "the path forward isn't necessarily writing new kernels from scratch — it's pulling in the existing upstream performance PRs and configs." This is the strategic pivot that defines the message.

Phase 7: Expectation Setting. The assistant calculates the realistic ceiling (40-100 tok/s) and explicitly states that 1000 tok/s is not achievable on this hardware. This is a crucial service to the user—without this reality check, the team could invest weeks of effort chasing an impossible target.

Phase 8: Verification. The assistant executes two commands to verify the hypotheses: a git log check (confirming missing commits) and a device query (confirming hardware specs). This grounds the reasoning in empirical evidence.

Phase 9: Synthesis. The assistant summarizes the findings and commits to presenting a plan, having gathered the necessary information.

This thinking process exemplifies several hallmarks of expert engineering reasoning:

The Broader Significance

This message matters beyond its immediate context because it illustrates a universal pattern in engineering: the moment when local investigation meets community knowledge. The assistant had spent hours profiling, tuning, and diagnosing—work that was necessary but insufficient. The upstream research provided the missing piece: evidence that others had already solved parts of the problem, and that the local deployment was missing those solutions.

The message also demonstrates the importance of maintaining a critical perspective on one's own data. The assistant could have confidently asserted that sparse decode was the bottleneck (63% is hard to argue with) and proceeded with the split-K rewrite. But the willingness to question that conclusion—to notice that upstream data told a different story and to investigate the discrepancy—prevented a potentially wasted effort. The split-K rewrite would have helped, but it would have addressed only part of the problem while leaving the FP8 GEMM autotune gap unaddressed.

Finally, the message models how to deliver bad news effectively. The assistant doesn't simply tell the user that 1000 tok/s is impossible—it provides the evidence, the reasoning, and a concrete alternative path (40-100 tok/s through upstream integration). This transforms a disappointing reality into an actionable plan.

Conclusion

Message 12451 is a masterclass in engineering decision-making under uncertainty. It captures the moment when a deep-dive investigation converges with external research to produce a fundamentally better strategy. The assistant's willingness to question its own profiling data, to seek out community knowledge, and to pivot from a custom-kernel approach to an upstream-integration approach demonstrates the intellectual flexibility that distinguishes exceptional engineers.

The message also serves as a cautionary tale about the dangers of working in isolation. Without the upstream research, the team might have invested weeks in a split-K kernel rewrite that would have delivered at most 2-3× improvement, missing the 4-5× gain available from FP8 GEMM autotuning alone. The lesson is clear: before building, first look to see what others have already built.

For the reader who has followed this session from the beginning, message 12451 is the turning point—the moment when the optimization effort shifts from "how do we make this work" to "how do we integrate what already works." It is a reminder that in the fast-moving world of ML infrastructure, standing on the shoulders of giants is not just efficient—it is essential.