The Turning Point: How Split-K Parallelization Unlocked 4× Latency Improvement and Redefined the Optimization Frontier

Introduction

In the high-stakes world of large language model inference on cutting-edge hardware, optimization is rarely a straight line. It is a cycle of measurement, hypothesis, implementation, and re-measurement—each iteration peeling back a layer of bottleneck to reveal the next one beneath. Few messages in a coding session capture this iterative drama as cleanly as message 12570 in the DeepSeek-V4-Flash deployment on 8× NVIDIA RTX PRO 6000 Blackwell GPUs (sm_120 architecture). This message is a hinge point: it simultaneously celebrates a major victory, performs a critical diagnostic, and sets the stage for the next phase of work. It is the moment the assistant confirms that a carefully engineered split-K parallelization strategy has more than doubled throughput across the board—and then immediately turns to profiling to understand what remains.

To appreciate why this message matters, one must understand the problem it solved. The assistant had previously built a custom MMA (matrix-matrix assembly) sparse-MLA decode kernel using Triton's tl.dot tensor-core operations, replacing a per-head SIMT kernel that was re-reading the KV cache 64× redundantly. That initial MMA kernel delivered a substantial improvement at high concurrency—C=64 throughput roughly doubled from 29.71 to 59.56 tokens/second—but it came with a troubling regression: at C=1 (single concurrent request), throughput actually dropped from 11.47 to 8.13 tok/s, and per-token latency ballooned to 122 milliseconds. The root cause was occupancy: with only two thread blocks for the entire GPU (188 SMs), the kernel left the vast majority of compute units idle. The user explicitly chose to address this before moving on to other optimizations, and the assistant responded by designing a split-K parallelization scheme that partitions the topk attention computation across multiple thread blocks and merges the results with a log-sum-exp (LSE) combine kernel. Message 12570 is the moment that design meets reality.

The Message: A Victory Lap and a Diagnostic Pivot

The message opens with the assistant's reasoning trace, which immediately signals the tone: measured satisfaction mixed with forward-looking analysis.

Looking at the performance metrics, split-K is delivering remarkable improvements, particularly at C=1 where it more than doubled throughput from 8.13 to 22.43 compared to the non-split MMA version. The gains hold across all tested configurations, with C=16 and C=64 also showing substantial speedups.

The language is precise: "remarkable improvements" is not hyperbole when the numbers back it up. The assistant then walks through the throughput calculations, noting that at C=1 the per-token time dropped to 29.78 milliseconds, translating to 33.6 tokens/second in steady generation—a 4× improvement over the 122 ms per-token latency of the MMA-only kernel. The reasoning reveals a sophisticated understanding of how concurrency affects measurement: "the math is getting trickier since I need to account for how TPOT relates to actual step time when multiple requests are advancing in parallel." This is not a superficial analysis; the assistant is grappling with the genuine complexity of benchmarking inference servers where batch processing interleaves multiple requests.

The key analytical insight comes next:

Looking at the marginal cost per additional request, it's hovering around 14-15ms across different batch sizes, which is similar to what we saw without the split-K optimization—the key insight being that split-K solved the occupancy problem at low concurrency (C=1 went from 122ms to 30ms) but the per-request overhead at high batch is still bottlenecked by the glue code and FP32 GEMM operations.

This is the essence of the message: split-K fixed the occupancy problem, but it did not change the per-request overhead. The asymptotic throughput ceiling is around 70 tokens/second, and C=64 is already close to it at 61.98 tok/s. The assistant correctly identifies that to push further, it must tackle the "glue code and FP32 GEMM operations"—the remaining bottlenecks that the profiling will soon reveal in stark detail.

The Benchmark Results: Quantifying the Victory

The assistant presents a clean comparison table that tells the story at a glance:

| C | old SIMT | MMA only | MMA + split-K | vs old | |---|---|---|---|---| | 1 | 11.47 | 8.13 | 22.43 | +96% (TPOT 122→30 ms, 4×) | | 16 | 26.61 | 40.97 | 56.46 | +112% | | 64 | 29.71 | 59.56 | 61.98 | +109% |

Every concurrency level roughly doubled compared to the original SIMT kernel. The C=1 case, which had been a regression, is now a 4× latency improvement. The split-K strategy worked exactly as designed: it recovered the low-occupancy regime by distributing the topk computation across more thread blocks, keeping the GPU's SMs busy even with a single request.

But the table also reveals something subtler. Compare the MMA-only column to the MMA+split-K column: at C=64, the improvement is modest (59.56 → 61.98, about 4%). The split-K optimization primarily helps at low concurrency. At high concurrency, the batch size itself provides enough work to keep SMs occupied, so the benefit of further parallelization is marginal. This is exactly what the assistant's reasoning predicted, and it validates the design choice to make NSPLIT adaptive—smaller splits when batch size is large, larger splits when batch size is small.

The Profile Pivot: Identifying the Next Bottleneck

Having confirmed the split-K victory, the assistant does not rest. It immediately launches a profiling run at C=32 to "confirm the attention share and size up the next target." This is the hallmark of systematic optimization: never declare victory without understanding what remains.

The profile results are revealing:

total GPU-kernel us: 5240525.8

 19.2%   1004.64ms  n=  231  void cutlass::Kernel2<cutlass_80_simt_sgemm_64x128_8x5_tn_align1>
 17.6%    923.16ms  n= 2541  void at::native::unrolled_elementwise_kernel<at::native::direct_copy_kernel_cuda>
 13.0%    682.07ms  n= 2783  void at::native::elementwise_kernel<128, 2, at::native::gpu_kernel_impl_nocast>
 12.8%   ...

The top entry is a cutlass_80_simt_sgemm kernel—a CUDA-core single-precision GEMM running in SIMT mode, consuming 19.2% of GPU time. This is the FP32 GEMM bottleneck the assistant suspected. The next two entries are elementwise copy kernels and generic elementwise kernels, consuming 17.6% and 13.0% respectively. Together, these represent the "glue code"—the PyTorch operations that shuffle data between the attention kernel and the MoE layers.

What is striking is what is not in the profile: the attention kernel itself. The custom MMA sparse-MLA decode kernel with split-K has been so thoroughly optimized that it no longer appears among the top GPU time consumers. The assistant's earlier work—replacing the per-head SIMT kernel, adding tensor-core MMA operations, implementing split-K parallelization—has successfully pushed attention from being the dominant cost (~57% in the original kernel) to being a minor contributor. The bottlenecks have shifted to the surrounding infrastructure.

Assumptions and Decisions Embedded in the Message

This message makes several implicit assumptions that are worth examining. First, the assistant assumes that the profiling at C=32 is representative of the bottleneck structure across all concurrency levels. This is a reasonable assumption given that the per-request overhead (the "marginal cost" of ~14-15ms) is similar across batch sizes, but it is an assumption nonetheless—the profile at C=1 might look different because the GPU is less occupied.

Second, the assistant assumes that the FP32 GEMM and elementwise operations are "glue code" that can be optimized independently. The cutlass SIMT sgemm kernel suggests that somewhere in the model, a matrix multiplication is falling back to single-precision CUDA cores instead of using tensor-core bf16 operations. The assistant had previously identified and flipped the forced-FP32 indexer bmm and MHC-pre linear to bf16, but this profile suggests there may be additional FP32 operations that were missed, or that the fix did not fully take effect.

Third, the assistant assumes that the asymptotic throughput ceiling of ~70 tok/s is a real bound imposed by these bottlenecks, not an artifact of the measurement methodology. This is a defensible assumption—the profile shows concrete GPU kernel times that account for the wall-clock time—but it leaves open the possibility that other factors (CPU-side scheduling, network I/O, NCCL communication) could become limiting once the GPU kernels are optimized further.

The Thinking Process: What the Reasoning Reveals

The agent reasoning in this message is unusually rich and reveals several layers of cognitive work. The assistant is simultaneously:

  1. Interpreting empirical results: Reading the benchmark numbers and translating them into physical insight about occupancy, marginal cost, and asymptotic limits.
  2. Validating a design hypothesis: The split-K strategy was explicitly designed to fix the C=1 occupancy problem, and the results confirm it worked. The assistant connects the design intent to the measured outcome.
  3. Calibrating expectations: The assistant notes that "the asymptotic throughput ceiling is around 70 tokens/second, and we're already close to it at C=64." This is a subtle but important move—it prevents the user from expecting unbounded gains from further optimization and frames the remaining work as diminishing returns.
  4. Planning the next iteration: The decision to profile at C=32 is strategic. C=32 is between the low-concurrency regime (where split-K matters most) and the high-concurrency regime (where the asymptotic ceiling dominates). It is a compromise that should reveal both the attention kernel's remaining footprint and the glue code's cost.
  5. Communicating results to the user: The table format, the clear "vs old" column, and the explicit mention of the 4× latency improvement at C=1 are all designed to make the results immediately understandable. The assistant is not just thinking—it is preparing to present findings.

Input Knowledge Required to Understand This Message

To fully grasp this message, one needs knowledge of several domains:

Output Knowledge Created by This Message

This message produces several distinct forms of knowledge:

  1. Quantitative validation: The benchmark table provides definitive evidence that split-K parallelization works as designed, with specific numbers for each concurrency level.
  2. Bottleneck diagnosis: The profile at C=32 identifies the next targets: FP32 GEMM (19.2%) and elementwise/copy operations (30.6% combined). This knowledge directly informs the next optimization phase.
  3. Asymptotic bound estimation: The assistant's calculation of a ~70 tok/s ceiling gives the user a realistic expectation of what is achievable without addressing the glue code.
  4. Design principle: The insight that "split-K solved the occupancy problem at low concurrency but the per-request overhead at high batch is still bottlenecked by the glue code" is a generalizable lesson about where different optimization techniques apply.
  5. Methodology: The approach of profiling at an intermediate concurrency level (C=32) to get a representative picture of bottlenecks is itself a methodological contribution.

Mistakes and Incorrect Assumptions

Were there any mistakes in this message? The assistant's analysis appears sound, but there are a few potential issues worth noting:

Significance in the Broader Campaign

Message 12570 sits at a critical juncture in the optimization campaign. The split-K implementation (messages 12562–12567) was the last major kernel-level change. After this message, the assistant will pivot to the "glue code" and discover something unexpected: the elementwise operations are not generic overhead but the DSA indexer torch fallback computing scores over the full 1M-token max context every decode step. This discovery will lead to a 17.9× throughput breakthrough at C=64, dwarfing everything that came before.

In retrospect, message 12570 is the calm before the storm. The assistant believes it has thoroughly optimized the attention kernel and is now turning to "mopping up" the glue code. It does not yet know that the glue code contains a hidden monster—the O(max_context) indexer—that will transform the optimization landscape. The profile shown in this message, with its 19.2% FP32 GEMM and 30.6% elementwise operations, is actually a misdirection: those numbers are symptoms of the indexer problem, not independent bottlenecks.

But this is not a failure of the analysis. The assistant's methodology is sound: measure, hypothesize, implement, re-measure. The split-K optimization was necessary—it fixed the C=1 regression and improved throughput across the board. Without it, the subsequent indexer fix would have been less impactful because the attention kernel itself would have remained a bottleneck. Each optimization layer builds on the previous one.

Conclusion

Message 12570 is a masterclass in systematic optimization. It demonstrates how to validate a complex kernel change with rigorous benchmarking, how to interpret results to extract physical insight, how to communicate findings clearly to a collaborator, and how to immediately pivot to the next diagnostic step. The assistant's reasoning reveals a deep understanding of GPU architecture, attention mechanisms, and inference serving dynamics.

The split-K parallelization strategy, implemented across several preceding messages, is validated as a complete success: it recovers the C=1 regression, doubles throughput across all concurrency levels, and delivers a 4× latency improvement for single requests. But the assistant does not rest on this victory. It immediately profiles to understand what remains, identifying the FP32 GEMM and elementwise operations as the next targets.

What the assistant does not yet know is that these targets are not independent bottlenecks but symptoms of a deeper problem—the indexer computing over the full max context every decode step. That discovery will come in the next phase and will deliver a 17.9× improvement that makes the gains in this message look modest by comparison. But that does not diminish the importance of this moment. The split-K optimization was necessary groundwork, and the profiling methodology established here will prove essential for the indexer diagnosis.

In the end, message 12570 is about the discipline of optimization: celebrate the win, but never stop measuring. The next bottleneck is always waiting.