Documenting the Ceiling: How a Single SUMMARY.md Captured the sm_120 Bottleneck on Blackwell

In the middle of an intensive optimization campaign to deploy DeepSeek-V4-Flash on eight NVIDIA RTX PRO 6000 (Blackwell sm_120) GPUs, the assistant paused to write a file. That file—/root/dsv4_bench/SUMMARY.md—was not a deliverable in itself, but it crystallized the entire arc of the session: the architecture, the benchmarks, the incremental gains, and the hard ceiling that no amount of configuration tuning could break through. Message [msg 12393] is the moment when the assistant transitions from active experimentation to consolidation, transforming a sequence of failed optimizations and partial victories into a coherent diagnosis of the fundamental bottleneck.

This article examines that single message in depth: why it was written, what knowledge it encodes, the assumptions it makes, the thinking process that shaped it, and the mistakes embedded in its conclusions. It is a case study in how a well-crafted summary can be as valuable as the experiments it documents—and how the act of writing forces clarity about what has been learned.

The Context: An Optimization Campaign That Hit a Wall

To understand message [msg 12393], one must first understand what preceded it. The assistant had been tasked with deploying DeepSeek-V4-Flash—a 284-billion-parameter Mixture-of-Experts model with 13 billion active parameters per token—on a machine with eight RTX PRO 6000 GPUs using SGLang. The headline technical challenge was that these GPUs use the sm_120 architecture (Blackwell), which lacks native support for the high-performance CUDA kernels that SGLang typically relies on: DeepGEMM, trtllm-gen, and the FP4 C4 indexer are all architecture-gated to sm_100 (Hopper) and simply do not run on Blackwell.

The optimization campaign that preceded this message (detailed in Chunk 1 of Segment 67) was methodical and data-driven. The assistant pulled the latest upstream SGLang, generated SM120 FP8 GEMM autotune configurations, and enabled MTP/EAGLE speculative decoding. Each lever was benchmarked individually. FP8 configs gave only ~6% improvement (because FP8 GEMM constitutes only 6% of decode time). NCCL LL had negligible effect (communication is only 2% of the profile). MTP boosted single-request throughput by 47% but provided zero gain at concurrency due to verifier saturation. A profile at concurrency 16 revealed the decisive bottleneck: the MoE slot-GEMV and sparse-decode attention kernels consumed 39% and 38% of GPU time respectively, both running on CUDA cores (SIMT) rather than the sm_120a FP4/FP8 tensor cores. This explained the user's observation of less than 1% tensor-pipe utilization.

The assistant then pivoted to the highest-leverage fix: switching from the stock MXFP4 checkpoint to the official NVIDIA NVFP4 quantization, which routes MoE execution through tensor-core paths (Marlin W4A16 or native cutlass FP4 grouped GEMM). PR #25820 was fetched and applied for correct NVFP4 auto-detection, and a 149 GB checkpoint was downloaded. Both NVFP4 backends delivered identical throughput (~28 tok/s at C=16, a ~24% improvement over baseline), confirming that MoE was now on tensor cores but attention remained the dominant bottleneck.

This is where message [msg 12393] arrives. The assistant has exhausted every configurable lever. The NVFP4 checkpoint has been deployed. The PD disaggregation pipeline—prefill on GPU0-3/NUMA0, decode on GPU4-7/NUMA1, KV transfer over NIXL/UCX—is verified working end-to-end. The benchmarks have been collected. Now it is time to write it all down.## The Structure of the Summary: What Gets Documented and Why

The SUMMARY.md file that the assistant writes in message [msg 12393] is structured in five sections: Stack, Single-node TP4, PD Disaggregation, Finding, and Launch Scripts. Each section serves a distinct epistemic purpose.

The Stack section is a precise inventory of every software component in the deployment pipeline. It lists the model (DeepSeek-V4-Flash, 284B/13B-active, 46 shards, ~146 GB), the SGLang commit hash (735a256), the virtual environment (venv_sglang211 with torch 2.11+cu130), every dependency with its version (flashinfer 0.6.12, sglang-kernel 0.4.3, tilelang 0.1.8, tokenspeed_mla 0.1.6, transformers 5.8.1, NIXL 1.3.0 + UCX, sglang-router 0.3.2), and the critical annotation about which paths are available on sm_120 versus which are architecture-gated off. This is not merely documentation—it is a reproducibility contract. If someone returns to this deployment in six months, they need to know exactly which commit of SGLang was used, because the difference between 735a256 and a later commit could be the difference between working NVFP4 detection and a silent fallback to MXFP4.

The Single-node TP4 section records the baseline: 10.2 tok/s at batch size 1, ~24 tok/s aggregate at batch size 8, with a TPOT of 94.1ms and TTFT of 696ms. These numbers are the control condition for every subsequent experiment.

The PD Disaggregation section records the headline deliverable—prefill-decode disaggregation works end-to-end—and then gives the hard numbers at three concurrency levels. The pattern is telling: at c=1, throughput is 9.95 tok/s (essentially identical to single-node). At c=8, it rises to 24.5 tok/s. At c=16, it plateaus at 25.0 tok/s. The time-to-first-token degrades from 893ms at c=1 to 8789ms at c=16, but aggregate throughput barely budges beyond c=8. This saturation pattern is the fingerprint of a decode-bound system.

The Finding section is the analytical core. It states bluntly: "Throughput is decode-bound by the sm_120 fallback kernels." It identifies the specific architectural gate: "the SM100 fast paths (DeepGEMM, trtllm-gen mxfp4, FP4 C4 indexer) are arch-gated off." And it delivers the crucial insight: "PD (isolating decode from prefill) does not raise aggregate tok/s here vs single-node—decode is the bottleneck, not prefill interference or KV transfer."

This last point is worth dwelling on. Prefill-decode disaggregation is a technique designed to solve a specific problem: when prefill and decode compete for the same GPU resources, prefill (which is compute-intensive and memory-bandwidth-heavy) can starve decode (which is latency-sensitive and memory-bound). By separating them onto different GPUs, PD should improve decode throughput. But it only helps if decode is not already the bottleneck. In this deployment, decode is so slow—running on fallback kernels that use CUDA cores instead of tensor cores—that isolating it from prefill interference makes no measurable difference. The bottleneck has simply moved from resource contention to kernel efficiency.

The Thinking Process: What the Assistant's Reasoning Reveals

The assistant's reasoning block at the top of message [msg 12393] is short but revealing. It confirms the deployment state—prefill on NUMA0 CPUs 0-31, decode on NUMA1 CPUs 32-63, all three components healthy—and then states the intention: "Now I'll write a deployment summary file to document the benchmarks, configuration, and findings."

The brevity is itself significant. Earlier messages in this session contained extensive reasoning about which optimization lever to try next, how to interpret ambiguous benchmark results, and what the next experiment should be. By message [msg 12393], the assistant has reached a terminal state. There are no more levers to pull. The reasoning is not "what should I try next?" but "what have I learned?" This is the cognitive shift from exploration to consolidation.

The assistant also takes care to verify the NUMA binding explicitly, checking /proc/$pid/status for Cpus_allowed_list on both the prefill and decode processes. This is a double-check that the numactl pinning worked correctly—prefill on CPUs 0-31 (NUMA0) and decode on CPUs 32-63 (NUMA1). The confirmation that "comms stay on 1 CPU each" references the user's earlier requirement that communication overhead should not spread across all CPUs. This attention to process-level verification shows a systematic approach: the assistant does not assume the configuration worked; it checks.

Assumptions Embedded in the Summary

Every summary rests on assumptions, and this one is no exception. The most important assumption is that the benchmark methodology—random 256/256 prompts, three concurrency levels (1, 8, 16), the specific SGLang benchmarking tool—produces results that generalize to real workloads. A 256-token input and 256-token output is a reasonable proxy for many chat interactions, but it does not capture the behavior of long-context applications (e.g., document analysis with 32K-token inputs) or streaming use cases where output length varies widely.

The summary also assumes that the bottleneck diagnosis is complete. The finding attributes the throughput ceiling to sm_120 fallback kernels, and the evidence supports this: GPU profiling showed the _tiled_sparse_decode_kernel consuming 63% of decode time, launching only 64 blocks on ~170 SMs. But the summary does not consider the possibility of compounding bottlenecks—for example, that the NIXL/UCX KV transfer over PCIe (without RDMA or NVLink) might become a bottleneck at higher throughputs, even if it is invisible at the current 25 tok/s. This is not a mistake, but it is a limitation of the diagnosis: the summary identifies the current bottleneck, not the next bottleneck that would appear if the current one were removed.

Another assumption is that the NVFP4 quantization is correctly handling all expert computations. The summary notes that both NVFP4 backends (Marlin W4A16 and native cutlass FP4 grouped GEMM) delivered identical throughput, which is strong evidence that MoE is now on tensor cores. But the summary does not verify that the NVFP4 checkpoint produces identical output quality to the stock MXFP4 checkpoint. The assumption is that NVIDIA's NVFP4 format is a faithful re-quantization, but the summary does not include a quality comparison (e.g., perplexity on a held-out dataset or a human evaluation of sample outputs).

Mistakes and Incorrect Assumptions

While the summary is largely accurate, there are subtle issues worth examining. The most significant is the framing of the throughput ceiling as an architectural limitation of sm_120. The summary states that "the SM100 fast paths... are arch-gated off," implying that the hardware itself cannot run these kernels. This is true in a narrow sense—the CUDA compute capability version (sm_120 vs sm_100) determines which instruction sets are available, and DeepGEMM uses sm_100-specific instructions. But the summary does not distinguish between cannot run and has not been ported. In principle, the DeepGEMM kernels could be rewritten for sm_120; the fact that they have not been is a software availability issue, not a hardware capability issue. The RTX PRO 6000 has tensor cores (sm_120a), and those tensor cores can execute FP4 and FP8 matrix operations—they just need kernels written for their instruction set. The summary's language ("arch-gated off") subtly shifts blame from the software ecosystem to the hardware, which may mislead a reader into thinking Blackwell is fundamentally incapable of fast inference.

The summary also makes a minor error in the throughput numbers. It reports "bs=8: ~24 tok/s aggregate" for single-node TP4, but the PD benchmark at c=8 reports 24.5 tok/s and at c=16 reports 25.0 tok/s. The near-identity of these numbers across vastly different concurrency levels (8 vs 16) and different GPU allocations (4 GPUs vs 8 GPUs) is presented as evidence of a decode bottleneck. But the summary does not note the statistical uncertainty in these measurements. A single benchmark run at each concurrency level, with a fixed prompt length and output length, does not provide confidence intervals. The difference between 24.5 and 25.0 tok/s is within the noise floor of any realistic benchmark, yet the summary treats it as a meaningful plateau.

Input Knowledge Required to Understand the Message

To fully grasp message [msg 12393], a reader needs substantial background knowledge. They need to understand what prefill-decode disaggregation is and why it is used—specifically, that separating prefill and decode onto different GPUs prevents prefill from starving decode of GPU resources. They need to know what NIXL and UCX are (NVIDIA's communication library and the Unified Communication X framework, respectively) and why KV transfer over PCIe without RDMA or NVLink is a potential bottleneck. They need to understand tensor parallelism (TP4) and how it distributes model layers across GPUs. They need to know what CUDA compute capability versions mean (sm_120 vs sm_100) and why certain kernels are architecture-gated. They need to understand the difference between CUDA cores and tensor cores, and why running matrix operations on CUDA cores is dramatically slower. They need to know what MXFP4 and NVFP4 quantization formats are, and why switching between them can route computation through different kernel paths.

The reader also needs to understand the significance of the specific numbers. A TPOT of 94ms means roughly 10.6 tokens per second, which is usable for interactive applications but far below the throughput needed for serving many concurrent users. The TTFT of 893ms at c=1 is acceptable, but the TTFT of 8789ms at c=16 is not—users would experience an 8-second delay before the first token. The aggregate throughput of 25 tok/s across eight GPUs means each GPU is producing roughly 3 tok/s, which is extremely low by modern standards (the user's roofline analysis suggested 300-600 tok/s should be possible).

Output Knowledge Created by This Message

Message [msg 12393] creates several distinct pieces of output knowledge. The most obvious is the SUMMARY.md file itself—a permanent record of the deployment configuration, benchmark results, and bottleneck diagnosis. This file serves as a reference for anyone who needs to reproduce or extend this work.

But the message also creates tacit knowledge that is not captured in the file. The assistant's decision to write the summary at this specific moment—after the NVFP4 deployment, after the PD benchmarks, after exhausting all configurable levers—implicitly communicates that the optimization campaign has reached a natural end. There is no "next step" section in the summary. There is no list of further experiments to try. The message says, in effect: we have done everything we can with configuration changes, and the remaining gap requires a fundamentally different approach (custom kernel development).

The message also creates knowledge about the relationship between PD disaggregation and kernel efficiency. Before this experiment, one might reasonably assume that PD disaggregation would improve throughput in any deployment where prefill and decode share GPUs. The message demonstrates that this assumption is false: PD only helps when decode is not already the bottleneck. This is a counterintuitive finding that has implications for how PD should be evaluated in future deployments.

The Broader Significance: When Documentation Becomes Diagnosis

Message [msg 12393] is, on its face, a mundane act of documentation. The assistant writes a summary file, verifies process placement, and reports results. But in the context of the full session, this message represents something more significant: the moment when the assistant transitions from an experimental mindset to a diagnostic one.

The optimization campaign that preceded this message was characterized by a relentless "try the next lever" approach. When FP8 configs gave only 6% improvement, the assistant tried NCCL LL tuning. When that failed, the assistant tried MTP speculative decoding. When MTP failed at concurrency, the assistant switched checkpoint formats to NVFP4. Each experiment was a hypothesis test: "will this lever improve throughput?" Each one returned a negative or marginal result.

Message [msg 12393] is where the assistant stops testing hypotheses and starts synthesizing them into a coherent diagnosis. The finding section of the SUMMARY.md is not a list of failed experiments—it is a causal model. It identifies the root cause (sm_120 fallback kernels), traces the mechanism (CUDA cores instead of tensor cores), explains why PD does not help (decode is the bottleneck, not prefill interference), and quantifies the gap (the SM100 fast paths are arch-gated off).

This is the most valuable kind of knowledge an optimization campaign can produce. A successful optimization—finding the right config flag that doubles throughput—is valuable but fragile; it may not transfer to a different model, different hardware, or different software version. A causal diagnosis of why throughput is limited is durable. It tells future engineers what they need to fix, not just what they need to set.

Conclusion: The Art of Knowing When to Stop

The most important decision in message [msg 12393] is not what the assistant wrote in the summary, but the implicit decision to stop optimizing and start documenting. The assistant could have continued trying more levers: different tensor parallelism strategies, different memory fraction settings, different CUDA graph configurations, different router policies. But the assistant recognized that the remaining gap—a factor of 10-20× below the user's throughput target—could not be closed by configuration changes alone. It required custom kernel development, a multi-week engineering effort akin to the earlier K2.6 work that delivered 3-6× gains through a custom verify attention kernel.

The SUMMARY.md file that the assistant writes in message [msg 12393] is, in this sense, a boundary object. It marks the boundary between what configuration can achieve and what requires kernel engineering. It tells the user: we have reached the ceiling of the current approach, and here is the data that proves it. This is not a failure—it is the most valuable output of the entire optimization campaign. Without this diagnosis, a future engineer might waste weeks trying the same levers that the assistant has already exhausted. With it, they know exactly where to invest their effort: in writing sm_120 tensor-core kernels for sparse attention and MoE computation.

The message also demonstrates a crucial engineering discipline: the discipline of writing down what you have learned before moving on. It is tempting, after a long optimization campaign, to jump straight to the next task—to start building the custom kernel, to try a different model, to reconfigure the hardware. The assistant resists this temptation and takes the time to consolidate knowledge into a permanent record. This is the mark of a mature engineering process, and it is why message [msg 12393] deserves careful study.