The 0.7 tok/s Revelation: When Speculative Decoding Collapses to Autoregressive Floor

Introduction

In the high-stakes world of large language model inference, speculative decoding is supposed to be the magic bullet—a technique where a smaller, faster "drafter" model proposes multiple candidate tokens, and the large "target" model verifies them in parallel, yielding several tokens per step instead of one. But magic bullets sometimes misfire. In message [msg 12174] of a months-long coding session deploying and optimizing the Kimi K2.6 model on RTX PRO 6000 Blackwell GPUs, the assistant confronts a sobering reality: the DDTree speculative decoder, after weeks of custom CUDA kernel development, kernel tuning, CUDA graph integration, and KV cache defragmentation, is committing only one token per step on long-context workloads. The decode throughput at 185k tokens has collapsed to 0.7 tokens per second—essentially the raw autoregressive speed with no speculative benefit whatsoever.

This message is a moment of reckoning. It is not a message that builds something new; it is a message that interprets disappointing data and makes a critical diagnostic decision. The assistant must decide whether the problem is a bug in the drafter, a metrics reporting issue, a fundamental limitation of the approach, or simply an undertrained model. The reasoning in this message—and the assumptions it makes—shapes the entire subsequent trajectory of the project.

The Message in Full

The assistant writes:

## Agent Reasoning Looking at these metrics, DDTree is performing much worse than I expected at long contexts with the synthetic prompt — it's only committing a single token per step, with virtually no draft acceptance beyond that bonus token. This suggests spec decoding isn't providing the speedup I anticipated in this scenario.

>

The decode throughput values (118, 20, 6.3, 2.8, 1.4, 0.7 tok/s) are essentially measuring step latency since commit_len≈1, and the step time clearly grows with context length. At 1.4k context I'm seeing around 8ms per step, but at 185k it balloons to 1.43 seconds per step. The drafter isn't helping at all here — either it's undertrained for long contexts or the synthetic repetitive text is causing it to mispredict. But there's something odd: avg_accepted_depth=3 suggests the accepted path reached depth 3 in the tree, yet avg_commit_len=1 and avg_accepted_drafts=1 imply only a single token was actually committed. These metrics seem contradictory and I need to reconcile what they're actually measuring. So the decode performance at long context is essentially raw autoregressive speed with no speculative benefit — commit length of 1 means the drafter isn't predicting ahead. Looking at the raw step times, there's a clean linear relationship above 10k tokens: roughly 7.7ms per additional 1k of context, which confirms attention is the bottleneck. I should save these measurements to the results file with a note about the commit length caveat and then decide on next steps.

>

Key finding from the metrics: avg_commit_len=1.00 — DDTree is committing only ~1 token/step on this continuation workload, so these decode numbers reflect essentially the raw C=1 autoregressive floor (the drafter isn't helping here; acceptance collapses on novel prose, consistent with the known undertrained tmp-long drafter). That makes this a clean measurement of raw per-step decode latency vs context. Let me record the results:

>

[write] /home/theuser/glm-kimi-sm120-rtx6000bw/kdtree-engine/bench_results_ct200/context_scaling.txt Wrote file successfully.

The Context: What Led to This Moment

To understand why this message matters, we must understand the journey that preceded it. The assistant and user had been working for weeks—across multiple segments and chunks—to deploy the Kimi K2.6 model with DFlash speculative decoding on RTX PRO 6000 Blackwell GPUs. This was not a simple deployment. The Blackwell RTX PRO 6000 uses the sm_120 architecture, which is a consumer-grade variant of Blackwell that lacks the advanced tensor core instructions (wgmma, TMA, tcgen05) found in the datacenter B300/B200 GPUs. Every optimized MLA (Multi-head Latent Attention) kernel—FlashMLA, cutlass-MLA, flashinfer-MLA—was compiled only for sm_90a, sm_100a, or sm_103a. None supported sm_120.

The assistant had to build a custom verify attention kernel from scratch, optimize it for sm_120's smaller 100KB shared memory, make it CUDA-graph-capture-safe, and integrate it into SGLang's live service. That kernel, after extensive tuning (increasing NSPLIT from 16 to 64, adding 128-bit vectorized bf16 KV loads), achieved a dramatic 3–6× end-to-end decode speedup over the Triton baseline. KV cache defragmentation was implemented. The service was extended to 200k context length. The bottleneck had been cleanly identified as MoE (Mixture-of-Experts) imbalance at batch size 1.

And then the assistant ran the context-length scaling benchmark—and discovered that the speculative decoder itself wasn't working.

The Metrics Puzzle: A Contradiction That Demands Resolution

The most intellectually interesting aspect of this message is the metrics contradiction the assistant identifies. The DDTree metrics report:

• avg_actual_nodes=9.00 — the tree contains 9 nodes (the drafter's candidate tree)
• avg_accepted_drafts=1.00 — only 1 draft token was accepted per step
• avg_commit_len=1.00 — only 1 token was committed per step
• avg_accepted_depth=3.00 — the accepted path reached depth 3 in the tree These numbers are contradictory on their face. If the accepted path reached depth 3 in the tree, how can only 1 token be committed? The assistant flags this explicitly: "These metrics seem contradictory and I need to reconcile what they're actually measuring." This is a crucial moment of intellectual honesty. The assistant does not simply accept the metrics at face value or jump to a conclusion. It acknowledges the contradiction and recognizes that the metrics need reconciliation. This kind of disciplined reasoning—pausing to question data that doesn't make sense—is the hallmark of a rigorous engineer. The assistant's implicit resolution is that avg_accepted_depth may measure the depth of the path selected within the tree (the branch chosen for verification), while avg_commit_len measures how many tokens actually pass the target model's verification. The drafter proposes a tree; the target model verifies it; but if the target model's verification rejects most of the draft tokens, only a short commit length results even though the drafter's tree was deep. This is consistent with an undertrained drafter that proposes plausible-looking but ultimately incorrect continuations.

Assumptions Made

The assistant makes several key assumptions in this message:

1. The drafter is undertrained. The assistant references the "known undertrained tmp-long drafter" as the likely cause of the collapse. This assumption is based on prior knowledge from the project's history—the drafter was a temporary/long-context variant that had not been fully trained. This is a reasonable inference, but it's also a hypothesis that could be tested (e.g., by trying a different drafter or evaluating on simpler text).

2. The synthetic repetitive text is not the cause. The assistant considers but ultimately dismisses the possibility that the synthetic benchmark text (generated by repeating a pattern) is causing the drafter to mispredict. The reasoning is that if the text is repetitive, the drafter should actually benefit from the predictability—repetitive text should be easier to predict, not harder. The assistant correctly leans toward the undertrained-drafter explanation.

3. The metrics are accurate. The assistant trusts that avg_commit_len=1.00 genuinely reflects the system's behavior and is not a measurement artifact. Given the earlier contradiction with avg_accepted_depth, this trust is conditional—the assistant acknowledges the puzzle but proceeds with the commit_len as the primary signal.

4. The step time scaling is purely attention-bound. The assistant observes a clean linear relationship (7.7ms per additional 1k of context above 10k tokens) and attributes this entirely to attention computation. This is a sound inference given that the verify attention kernel's cost scales linearly with KV cache size, but it assumes no other bottlenecks (e.g., memory bandwidth, CPU orchestration) are contributing at long context.

5. The benchmark data is still valuable. Despite the speculative decoder not working, the assistant decides that the measurements represent a "clean measurement of raw per-step decode latency vs context." This is a pragmatic reframing: the data is not wasted; it characterizes the system's fundamental attention scaling even if the speculative benefit is absent.

Input Knowledge Required

To fully understand this message, one needs:

• Understanding of speculative decoding: How a drafter proposes candidates and a target model verifies them, with the goal of committing multiple tokens per step.
• DDTree specifics: The tree-based speculative decoding approach where the drafter builds a tree of candidate continuations rather than a single chain.
• The project's history: The tmp-long drafter was known to be undertrained from earlier work. The custom verify kernel, CUDA graphs, and defragmentation were all built in preceding chunks.
• SGLang's metrics system: The meaning of avg_commit_len, avg_accepted_drafts, avg_accepted_depth, and avg_actual_nodes in the DDTree metrics output.
• Attention scaling laws: The linear relationship between KV cache size and per-step decode latency in transformer attention.
• The benchmark methodology: Streaming single-prefill requests with timing at TTFT (time to first token) and decode rate from subsequent tokens.

Output Knowledge Created

This message produces several important outputs:

1. A documented finding: The assistant writes the results to bench_results_ct200/context_scaling.txt, creating a permanent record of the context-length scaling measurements with the caveat that they represent the autoregressive floor rather than speculative-decoding performance.

2. A diagnostic conclusion: The commit_len=1 finding is the key output—it tells the team that DDTree is not accelerating decoding on this workload with the current drafter. This is a critical input for deciding next steps (train a better drafter, use a different speculative decoding approach, or accept the autoregressive performance).

3. A reframed interpretation: The assistant transforms what could be seen as a failure (speculative decoding not working) into a useful measurement (clean attention scaling data). This is an important intellectual contribution—it prevents the benchmark effort from being wasted and provides a baseline for future improvements.

4. A metrics puzzle for future investigation: The contradiction between avg_accepted_depth=3 and avg_commit_len=1 is flagged but not resolved. This creates a todo item for deeper metrics investigation.

Mistakes and Incorrect Assumptions

The assistant's reasoning is generally sound, but there are potential issues:

The drafter undertraining hypothesis is plausible but unverified. The assistant does not run a diagnostic to confirm this—for example, testing the drafter on simpler text (e.g., repetitive patterns that should be easy to predict) or comparing against a known-good drafter. The conclusion is based on prior knowledge and inference, not direct evidence.

The metrics contradiction is not fully resolved. The assistant acknowledges the puzzle but does not dig into the SGLang source code to understand exactly what each metric measures. The resolution is implicit and somewhat unsatisfying. A more thorough investigation might have revealed a metrics bug or a subtlety in how DDTree reports its statistics.

The assumption that step time is purely attention-bound may miss other factors. At 1.43 seconds per step at 185k context, there could be contributions from CPU-side orchestration (tree building, mask computation), CUDA graph launch overhead, or memory bandwidth contention. The linear fit is suggestive but not definitive proof.

The assistant does not consider whether the synthetic prompt's repetitiveness could cause the drafter to overfit to a single pattern. If the prompt is highly repetitive (e.g., "The quick brown fox jumps over the lazy dog" repeated thousands of times), the drafter might predict the next repetition correctly but the target model might reject it as too obvious or incorrect in the latent space. This is a subtle failure mode that the assistant does not explore.

The Thinking Process: A Window into Engineering Judgment

The reasoning in this message reveals several aspects of the assistant's thinking style:

Pattern recognition: The assistant immediately recognizes the significance of avg_commit_len=1.00 and connects it to the decode throughput numbers. The pattern "step time grows with context, commit_len is 1, therefore decode tok/s = 1/step_time" is recognized and validated.

Hypothesis generation: The assistant generates two competing hypotheses for why the drafter isn't working: (1) the drafter is undertrained for long contexts, and (2) the synthetic text causes misprediction. Both are considered and weighed.

Contradiction detection: The assistant spots the inconsistency between avg_accepted_depth=3 and avg_commit_len=1 and flags it explicitly. This is a critical skill—many engineers would gloss over the contradiction or not notice it.

Pragmatic reframing: Rather than treating the failed speculative decoding as a disaster, the assistant reframes the data as a "clean measurement of raw per-step decode latency vs context." This turns a potential dead end into useful knowledge.

Disciplined documentation: The assistant saves the results to a file with appropriate caveats, ensuring the knowledge is preserved for future reference.

The Broader Significance

This message captures a moment that every engineer faces: the moment when a carefully built system fails to deliver its promised benefit. The custom CUDA kernel, the CUDA graph integration, the KV defragmentation—all of that work was in service of making speculative decoding fast. But if the speculative decoder itself doesn't propose acceptable tokens, the speed of verification is irrelevant.

The message is a reminder that in complex systems, bottlenecks can shift in unexpected ways. The team had been focused on making the verify kernel faster, assuming that the drafter would do its job. The benchmark revealed that the drafter was the weakest link all along—not because of kernel performance, but because of model quality.

This is also a lesson in the value of measurement. Without the comprehensive context-length benchmark, the team might have continued optimizing the verify kernel, achieving ever-faster step times while wondering why end-to-end throughput wasn't improving. The benchmark forced the real bottleneck into the open.

Conclusion

Message [msg 12174] is a masterclass in diagnostic reasoning under disappointing results. The assistant receives metrics that show speculative decoding has collapsed to the autoregressive floor, identifies a metrics contradiction, generates and weighs hypotheses, reframes the data as still valuable, and documents the findings. The reasoning is disciplined, intellectually honest, and pragmatic. The message does not solve the problem—it identifies it, which is the necessary first step. The subsequent trajectory of the project (training a better drafter, exploring alternative approaches) would be shaped by the conclusion reached in this message: that the drafter, not the kernel, is the bottleneck.