The 64-Block Barrier: Diagnosing a GPU Occupancy Bottleneck in DeepSeek-V4-Flash on Blackwell

In any high-performance computing optimization campaign, there comes a moment of clarity when speculation ends and measurement speaks. For the assistant working to deploy DeepSeek-V4-Flash on eight NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture), that moment arrived in message [msg 12445]. After weeks of configuration tuning, framework upgrades, and benchmark runs—each yielding incremental or negative gains—the assistant finally isolated the root cause of a stubborn throughput ceiling. The diagnosis was precise, quantified, and actionable: a single Triton kernel for sparse-MLA attention was launching only 64 GPU thread blocks on a machine with roughly 170 streaming multiprocessors, leaving nearly two-thirds of the hardware idle during decode. This article examines that message in depth: the reasoning that led to the discovery, the assumptions that preceded it, the output it produced, and the architectural lessons it reveals about deploying modern large language models on cutting-edge but architecture-gated hardware.

The Optimization Campaign That Led Nowhere

The context for this message is a sustained and methodical effort to achieve high-throughput inference for DeepSeek-V4-Flash, a mixture-of-experts (MoE) model with sparse attention, on Blackwell GPUs. The assistant had already deployed the model successfully using SGLang with prefill-decode disaggregation—prefill on one set of four GPUs and decode on another—and verified end-to-end correctness. But performance was far below expectations: roughly 10 tokens per second at batch size 1 and 23–25 tokens per second at concurrency 16, against a user target of approximately 1000 tokens per second.

The assistant had systematically exhausted every available configuration lever. NCCL communication tuning (LL/Ring protocols, channel counts) produced no measurable improvement—a profile later showed communication consumed only 2% of GPU time. CUDA graphs, which eliminate kernel launch overhead, were already enabled and contributed nothing to the raw compute bottleneck. The tilelang indexer fusion optimization failed to JIT-compile on the sm_120 architecture. Non-Marlin MoE backends were invalid for FP4-quantized experts. Expert parallelism (EP4) actually regressed throughput to 14 tokens per second due to PCIe all-to-all communication overhead at small batch sizes. Every knob had been turned, and none moved the needle.

This pattern of diminishing returns created the conditions for the message under analysis. The assistant needed a definitive, measurement-based answer to a question that configuration tuning could not resolve: where exactly is the time going?

The Smoking Gun: Grid = (B, H)

The answer came from a GPU trace captured with the PyTorch profiler across 40 decode steps. The trace, parsed in message [msg 12442], revealed a stark distribution: a single kernel named _tiled_sparse_decode_kernel consumed 2412 milliseconds out of 3830 total GPU kernel time—63% of all GPU compute. The next-largest contributor, the MXFP4 MoE slot-GEMV kernel, accounted for only 380 milliseconds (10%). NCCL all-reduce was a mere 85 milliseconds (2%). The bottleneck was not communication, not MoE computation, not the indexer or Sinkhorn iterations—it was the sparse-MLA decode attention kernel, and it was dominating the profile by a factor of six over anything else.

But identifying which kernel was slow was only half the answer. The assistant needed to understand why it was slow. That required reading the kernel source code, which the assistant did in message [msg 12444]. The critical line was found at line 246 of flash_mla_sm120_triton.py:

grid = (B, H)

At batch size 1, this produces a grid of exactly 64 blocks (1 batch × 64 query heads). On a GPU with approximately 170 streaming multiprocessors (SMs), that means only 38% occupancy—the majority of compute units sit idle while the 64 active blocks each serially iterate through all 512 top-k tokens in tiles of 16 or 32. The kernel has no parallelization across the top-k/KV dimension; each block marches through the full sequence of selected tokens one tile at a time, performing gather, matrix-vector multiply, and accumulation in a tight but fundamentally under-parallelized loop.

The assistant recognized this pattern immediately. It was the exact same pathology that had plagued the K2.6 verify attention kernel in earlier work—a low-occupancy, serial-over-KV decode kernel that the assistant had previously fixed with a split-K rewrite, achieving 3–6× speedup. The diagnosis in message [msg 12445] is not just a finding; it is a recognition of a recurring architectural pattern across different models and kernels on the same hardware.

The Split-K Prescription

With the root cause identified, the assistant articulated the fix with precision. The solution is a split-K strategy: reshaping the grid from (B, H) = (1, 64) to (B, H, NSPLIT) = (1, 64, 8) = 512 blocks, where each block processes a subset of the 512 top-k tokens. The partial results are then combined with an online softmax merge—a standard technique for parallelizing attention over the key-value dimension. This would achieve full occupancy on the ~170 SMs and potentially deliver a 3–6× speedup on the sparse-decode kernel alone.

The assistant's reasoning here is notable for its restraint. It does not claim that split-K will solve the entire throughput problem. Instead, it performs a careful ceiling analysis: even if the attention kernel were made instantaneous (zero milliseconds), the remaining kernels—MXFP4 MoE GEMV, FP8 matmul, elementwise operations—still consume 1418 milliseconds of GPU time. The maximum achievable speedup from optimizing attention alone is therefore 3830 / 1418 ≈ 2.7×. At the current 23 tokens per second at concurrency 16, that yields roughly 60 tokens per second. With a more realistic 5× attention speedup (not infinite), the total might reach 46 tokens per second. Even aggressive full-stack kernel work on sm_120 might push to 50–70 tokens per second.

This is the message's central intellectual contribution: a quantified, evidence-based ceiling that separates what is achievable through kernel optimization from what is fundamentally impossible on this hardware generation. The 1000 tokens per second target, the assistant concludes, likely comes from B200/SM100 benchmarks where the entire fused kernel stack—DeepGEMM, trtllm-gen, FP4 C4 indexer—is available and optimized. On sm_120, those fused kernels do not exist, and no amount of configuration tuning or single-kernel work can close the gap.

A Postmortem of Failed Levers

The message also serves as a retrospective explanation for why every earlier optimization attempt failed. The profile data makes this clear with brutal simplicity:

• NCCL tuning (LL/Ring protocols, channel counts) targeted communication, which is only 2% of GPU time. Even a 50% improvement in communication yields 1% overall gain—imperceptible.
• CUDA graphs eliminate kernel launch overhead, but the bottleneck is kernel execution time, not launch latency. The graphs were already doing everything they could.
• Tilelang indexer fusion failed to compile on sm_120/CUDA 13, and even if it had worked, it targets the indexer path, not the sparse-decode attention that dominates.
• Expert parallelism added PCIe all-to-all communication that outweighed any benefit from splitting experts across GPUs at small batch sizes. The pattern is consistent: every lever the assistant pulled was targeting the wrong bottleneck because the true bottleneck was invisible without a profile. This is a powerful lesson in performance engineering: configuration tuning without measurement is guesswork, and the most sophisticated optimization campaign can be rendered moot by a single kernel that the tuner never thought to examine.

The Output: PROFILE_FINDINGS.md

The message's concrete output is a file saved to the remote machine: /root/dsv4_bench/PROFILE_FINDINGS.md. This document captures the full analysis in a structured format:

• Total GPU kernel time: 3830 ms across the profiled window
• Kernel breakdown: the sparse-decode kernel at 63%, MoE GEMV at 10%, FP8 matmul at 6%, NCCL at 2%, and miscellaneous kernels filling the remainder
• Root cause: grid=(B,H) producing only 64 blocks at batch size 1, with each block serially iterating all 512 top-k tokens—the same low-occupancy pathology as the K2.6 verify kernel
• Why config levers failed: a concise explanation for each attempted optimization, grounded in the profile data
• Ceiling analysis: the mathematical limit of what single-kernel optimization can achieve (2.7× overall, ~60 tok/s at C=16)
• Real engineering lever: the split-K prescription, explicitly referencing the K2.6 precedent This document serves multiple purposes. It is a record of the investigation for the user, a reference for future optimization work, and a boundary object that separates what is known from what remains to be done. The assistant also verifies that the TP4 server is still running after the profiling session, ensuring that the diagnostic work did not disrupt the live service.

Architectural Lessons: sm_120 vs. SM100

Beneath the technical specifics, this message illuminates a broader architectural reality. The Blackwell GPU generation (sm_120) includes tensor cores capable of FP4 and FP8 matrix operations, but the high-performance fused kernel stack—DeepGEMM for MoE, trtllm-gen for attention, the FP4 C4 indexer—is compiled and optimized for SM100 (the B200 architecture). On sm_120, these operations fall back to Triton-based kernels or pure-PyTorch implementations that run on CUDA cores (SIMT) rather than tensor cores. The result is a dramatic performance gap that cannot be closed by configuration or even single-kernel rewrites; it requires porting the entire fused kernel stack to the new architecture, a multi-week engineering effort.

The assistant's earlier success with the K2.6 verify kernel—a custom split-K rewrite that achieved 3–6× speedup—demonstrates that targeted kernel work can yield substantial gains. But the ceiling analysis in this message shows that even a repeat of that success would leave the system at roughly 2–3× current throughput, still an order of magnitude below the target. The remaining gap is structural: the MoE GEMV kernel, the FP8 matmul, and the hundreds of small elementwise operations are all sm_120 fallbacks that collectively consume 37% of GPU time and cannot be eliminated without a full fused-kernel stack.

Conclusion

Message [msg 12445] represents the culmination of a rigorous, measurement-driven optimization campaign. The assistant moved from configuration tuning to profiling to kernel source analysis, and finally to a definitive, quantified diagnosis. The finding—that a single kernel launches only 64 blocks on a 170-SM GPU, achieving 38% occupancy—is specific, reproducible, and actionable. The split-K prescription offers a clear path to a 2–3× improvement. But the ceiling analysis is equally important: it sets realistic expectations, explains why previous optimizations failed, and identifies the architectural gating factor that separates what is achievable on sm_120 from what requires SM100.

For the user, this message provides clarity. The 1000 tokens per second target is not reachable on this hardware with the current software stack, but a bounded 2–3× improvement through split-K kernel work is feasible. For the assistant, it demonstrates the value of measurement over speculation, the power of recognizing recurring patterns across different models, and the discipline of quantifying ceilings rather than promising unbounded gains. In the broader narrative of the coding session, this message is the turning point where optimization ends and engineering begins—where the assistant stops tuning knobs and starts planning the kernel rewrite that will define the next phase of work.