The Honest Ceiling: When Prefill-Decode Disaggregation Meets sm_120 Fallback Kernels

Introduction

In the high-stakes world of large language model deployment on cutting-edge hardware, success is rarely binary. A model can load, generate correct output, and route requests through an elaborate multi-GPU orchestration system—yet still fall orders of magnitude short of production throughput targets. Message <msg id=12394> captures precisely this moment of reckoning. It is the closing summary of an intensive optimization campaign spanning multiple sessions, where an AI assistant deployed DeepSeek-V4-Flash—a 284-billion-parameter mixture-of-experts model with only 13 billion active parameters—across eight NVIDIA RTX PRO 6000 (Blackwell) GPUs using SGLang's prefill-decode (PD) disaggregation architecture.

The message is simultaneously a delivery report, a performance analysis, and a candid admission of a fundamental hardware bottleneck. It tells the story of a deployment that works correctly end-to-end but achieves roughly 25 tokens per second at concurrency—versus a user target of approximately 1,000 tok/s. The gap is not a configuration error or a software bug; it is a structural limitation imposed by the GPU architecture itself. The fast tensor-core paths for this model's FP4 quantized experts are arch-gated to the SM100 compute capability, while these Blackwell RTX PRO 6000 cards expose sm_120, forcing execution through slow fallback kernels that run on CUDA cores instead of tensor cores.

This article examines message <msg id=12394> in depth: why it was written, what decisions it reflects, the assumptions embedded within it, and the hard-won knowledge it creates. It is a case study in the difference between "works" and "works fast enough," and a reminder that even the most sophisticated software orchestration cannot overcome a hardware capability mismatch.

Why This Message Was Written

Message <msg id=12394> is the capstone of a long chain of work. To understand its motivation, one must trace back through the preceding segments. The assistant had been engaged in a multi-session effort to deploy and optimize large language models on Blackwell GPUs. Earlier segments covered building custom CUDA kernels for Kimi K2.6's speculative decoding engine ([msg 12385] context), benchmarking Qwen models, and ultimately pivoting to DeepSeek-V4-Flash—a model that promised higher performance through aggressive FP4 quantization.

The immediate predecessor messages show the assistant methodically building the deployment. In <msg id=12385>, launch scripts were written for the PD disaggregation setup: a prefill server on GPU0-3 pinned to NUMA0, and a decode server on GPU4-7 pinned to NUMA1, with KV cache transfer via NIXL/UCX. In <msg id=12386> and <msg id=12387>, both servers came online and the router was launched. In <msg id=12389> through <msg id=12392>, benchmarks were run at concurrency levels 1, 8, and 16. The results were sobering: throughput plateaued at approximately 25 tok/s regardless of concurrency.

Message <msg id=12394> is written to deliver these results to the user in a structured, honest format. It is not a triumphant "mission accomplished" message—it is a "here is what we achieved, here is where it falls short, and here is why" message. The assistant's motivation is transparent: to provide a complete picture of the deployment's status, to explain the root cause of the performance ceiling, and to offer the user a clear choice about where to invest effort next.

The structure of the message reflects this tripartite goal. It opens with "What was delivered" (three concrete achievements), presents "Benchmarks" (hard data), delivers a "Key finding (honest)" (the bottleneck diagnosis), clarifies the "nvfp4" situation (a common point of confusion), and closes with three actionable next steps. This is a status report designed to enable decision-making.

What Was Delivered: Three Achievements

The assistant lists three deliverables, each representing significant engineering effort:

1. Made DeepSeek-V4-Flash work on SGLang at sm_120. This required stopping the previous K2.6 service, freeing disk space, downloading the 146 GB FP4 checkpoint (46 shards), pulling a fresh SGLang main branch, and installing all CUDA 13 dependencies including flashinfer 0.6.12, sglang-kernel 0.4.3, tilelang, and NIXL. Notably, the assistant emphasizes that "no GEMM hand-porting was needed"—the SGLang fork already shipped in-tree sm_120 kernels, unlike the earlier K2.6 work which required building custom verify attention kernels from scratch.

2. Single-node TP4. The model loads and generates correctly on a single node with tensor parallelism across 4 GPUs. The output is coherent—the assistant tested with a question about PD disaggregation and got a sensible answer mentioning "Paris."

3. PD disaggregation. This is the headline architectural achievement: prefill running on GPU0-3 (NUMA0) with eager mode, decode running on GPU4-7 (NUMA1) with CUDA graphs, KV cache transfer via NIXL/UCX, and an SGLang router on port 8000 dispatching requests. CPU pinning was verified—prefill on CPUs 0-31, decode on CPUs 32-63. The entire pipeline works end-to-end correctly.

These three bullet points establish credibility. They show that the assistant did not fail to deploy the model—the deployment is complete and functional. The problem is not with the software stack but with the hardware's compute capability.

The Performance Reality

The benchmark table in the message tells a stark story:

| Config | tok/s | TPOT | TTFT | |---|---|---|---| | Single-node TP4 bs=1 | 10.2 | 94 ms | 696 ms | | PD c=1 | 9.95 | 94 ms | 893 ms | | PD c=8 | 24.5 | 282 ms | 2.0 s | | PD c=16 | 25.0 | 510 ms | 8.8 s |

Several observations stand out. First, single-node and PD achieve nearly identical throughput at the same batch size (approximately 10 tok/s at effective batch size 1). Second, PD scales with concurrency up to approximately 25 tok/s, then plateaus. Third, time-to-first-token (TTFT) degrades significantly with concurrency, from 893 ms at c=1 to 8.8 seconds at c=16—a 10× increase that indicates queue buildup on the decode side.

The assistant's interpretation is precise: "PD doesn't raise aggregate tok/s here—decode is the bottleneck, not prefill interference or KV transfer." This is a crucial insight. Prefill-decode disaggregation is designed to isolate the two phases so that prefill does not block decode (and vice versa). But if decode itself is the bottleneck—if the decode GPUs cannot generate tokens any faster regardless of how cleanly they are isolated—then PD offers no throughput benefit. The assistant is essentially dedicating 4 GPUs to decode and getting the same throughput as a single-node setup using 4 GPUs for everything.

The Honest Assessment: sm_120 Fallback Kernels

The core finding of message <msg id=12394> is the identification of the sm_120 fallback kernel bottleneck. The assistant explains that the auto-selected marlin→sm120_triton MXFP4 MoE path and the forced torch indexer-logits (via SGLANG_FP8_PAGED_MQA_LOGITS_TORCH) are running on CUDA cores rather than tensor cores. All the SM100 fast paths—DeepGEMM, trtllm-gen mxfp4, the FP4 C4 indexer—are "arch-gated off" on sm_120.

This is a hardware capability mismatch. The RTX PRO 6000 (Blackwell) cards have sm_120 compute capability, but the optimized tensor-core kernels for FP4 matrix multiplication and indexing were compiled only for sm_100 (the previous Blackwell variant). The model's FP4 experts, which should run on dedicated tensor-core hardware at high throughput, are instead being executed through Triton fallback paths that use general-purpose CUDA cores.

The assistant quantifies the gap: the card achieves approximately 10× lower throughput than vLLM's TP4 number for the same model. CUDA graphs are confirmed enabled; the GPUs are compute-bound in the fallback kernels. This is not a configuration issue that can be fixed by tweaking NCCL settings or changing the MoE backend—it is a fundamental software limitation that requires either new kernel implementations targeting sm_120 or a hardware upgrade to sm_100-capable cards.

The nvfp4 Clarification

A subtle but important point in the message is the clarification about "nvfp4." The stock checkpoint is already the FP4 path (expert_dtype: fp4), and SGLang routes FP4 experts through the Marlin/sm120_triton MXFP4 method on this GPU. The assistant notes that native nvfp4 cutlass kernels (nvfp4_scaled_mm_sm120.cuh exists) and the FP4 indexer are "not auto-selected on sm_120." This means the software stack has the capability to run FP4 on tensor cores for sm_120, but the automatic selection logic does not choose these paths—likely because they are not yet fully integrated or tested.

This clarification prevents a common misunderstanding: the user might think "nvfp4" is a different checkpoint format that would solve the problem. The assistant is saying no—the checkpoint is already FP4, and the issue is that the fast execution paths for that format are not activated on this hardware. The "actual optimization frontier" is a multi-day kernel effort to enable and tune these native sm_120 paths, similar in scope to the earlier K2.6 custom kernel work.

Assumptions and Required Knowledge

To fully understand message <msg id=12394>, the reader needs familiarity with several concepts:

Knowledge Created by This Message

Message <msg id=12394> creates several important pieces of knowledge:

  1. DeepSeek-V4-Flash can be deployed on sm_120 Blackwell GPUs with SGLang. The model loads, generates correct output, and works through PD disaggregation. This is non-trivial—the 146 GB checkpoint, the FP4 expert format, and the sm_120 kernel requirements all posed potential failure points.
  2. PD disaggregation works correctly on this hardware. The NIXL/UCX KV transfer, the NUMA pinning, the router orchestration—all function as designed. This is a valuable reference for anyone deploying PD on similar hardware.
  3. The sm_120 fallback kernel bottleneck is the primary performance limiter. The message provides a clear diagnosis: approximately 10× below expected throughput, with the bottleneck localized to the MoE and attention kernels running on CUDA cores instead of tensor cores.
  4. PD does not help when decode is the bottleneck. This is a counterintuitive finding—one might expect PD to always improve throughput by eliminating interference. The message demonstrates that PD's benefit is contingent on decode having compute headroom.
  5. The nvfp4 path exists but is not auto-selected. The hardware can potentially run FP4 on tensor cores, but the software stack does not choose these paths automatically. This identifies a concrete optimization target.
  6. A multi-day kernel effort is the next step. The assistant implicitly estimates the effort required to close the performance gap: "a multi-day kernel effort, same flavor as the K2.6 work."

The Thinking Process

The assistant's reasoning in message <msg id=12394> reveals several cognitive patterns:

Systematic enumeration. The message follows a clear logical structure: deliverables first, then data, then interpretation, then implications. This reflects the assistant's training to present information in a structured, digestible format.

Candid self-assessment. The "Key finding (honest)" section is notably frank. The assistant does not attempt to spin the results positively or blame external factors. It states plainly that PD "doesn't raise aggregate tok/s here" and that the ceiling is "the sm_120 decode kernels." This honesty is valuable for building trust and enabling correct decision-making.

Precision about what works vs. what doesn't. The assistant carefully distinguishes between the orchestration (which works correctly) and the kernel performance (which is inadequate). This prevents the user from conflating a deployment bug with a hardware limitation.

Forward-looking orientation. The message closes with three concrete next steps (MTP/EAGLE spec-decode, higher concurrency testing, kernel profiling), each representing a different strategy for attacking the bottleneck. This transforms the status report into a decision point.

Awareness of the user's mental model. The assistant anticipates the user might think "nvfp4" is a different format that would solve the problem, and proactively clarifies this. It also references the earlier K2.6 work as a point of comparison, assuming the user is familiar with that context.

Conclusion

Message <msg id=12394> is a masterclass in honest technical communication. It delivers a complex deployment status with precision, identifies the root cause of a severe performance shortfall, and provides actionable next steps—all while maintaining a clear distinction between what was achieved and what remains to be done.

The message's central lesson is that software orchestration, no matter how sophisticated, cannot compensate for a hardware capability mismatch. The PD disaggregation pipeline works flawlessly; the NIXL/UCX KV transfer functions correctly; the model generates coherent output. Yet the system achieves only 2.5% of its target throughput because the GPU's tensor cores cannot execute the model's FP4 experts at full speed.

For the reader, this message offers a template for reporting on ambitious technical projects: celebrate what works, quantify what doesn't, explain why, and lay out the path forward. In an era where AI infrastructure is pushed to its limits daily, such honest reckoning is more valuable than ever.