The 100× Gap: When Roofline Analysis Exposes a Structural Bottleneck in DeepSeek-V4-Flash Inference

Introduction

In the high-stakes world of large language model deployment, few moments are as revealing as the collision between theoretical hardware capability and measured performance. This article examines a single message (index 12515) from an opencode coding session where an AI assistant, after days of methodical optimization work deploying DeepSeek-V4-Flash on 8× NVIDIA RTX PRO 6000 Blackwell GPUs, confronts a devastating discrepancy: the model is delivering only 28 tokens per second at a concurrency of 16, when the hardware's memory bandwidth alone should support thousands. The message is a masterclass in diagnostic reasoning—a moment where the assistant abandons kernel-level micro-optimizations and performs a roofline analysis that reveals a ~100× gap between expectation and reality, forcing a fundamental re-examination of where the wall-clock time is actually going.

The Context: A Long Optimization Campaign

To understand the significance of this message, we must first appreciate the journey that led to it. The session (Segment 67) began with the deployment of DeepSeek-V4-Flash (a Mixture-of-Experts model with ~13B active parameters) on Blackwell GPUs using SGLang. The assistant had successfully achieved prefill-decode disaggregation—a sophisticated architecture where prefill runs on GPU 0–3 (NUMA 0) and decode runs on GPU 4–7 (NUMA 1), with KV cache transfer via NIXL/UCX and a router on port 8000. This was a genuine engineering achievement, verified end-to-end.

However, performance was abysmal: approximately 10 tok/s at batch size 1 and ~25 tok/s at C=16. The assistant systematically exhausted every configuration lever—NCCL LL+Ring tuning, CUDA graphs (already enabled), tilelang indexer fusion (which failed to JIT-compile on sm_120), non-marlin MoE backends (invalid for FP4 experts), and expert parallelism (which was worse due to PCIe all-to-all overhead). A GPU profile traced 63% of decode time to a single kernel: _tiled_sparse_decode_kernel, the sm_120 Triton fallback for sparse MLA attention, which launched only 64 blocks on ~170 SMs.

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

It was at this point that the user interjected with message 12514, a pointed challenge: "There is no reason why TP4 is so bad at C=16, at C=16 we should be in 300-600t/s territory at least. We have plenty of FLOPS, 64GB/s pcie on single CPU capable of P2P, and 1.9TB/s vram bw, we should be able to infer from those specs what to expect from this A10B model (=a lot)."

This pushback—asserting that the hardware should deliver an order of magnitude more throughput—is the catalyst for the target message. The assistant had been operating under the assumption that the ~28 tok/s was an unfortunate but real ceiling imposed by sm_120 fallback kernels. The user's challenge forces a complete re-evaluation.

The Target Message: A Roofline Awakening

Message 12515 opens with the assistant's agent reasoning, and the first sentence already shows a shift in perspective:

"At C=16 we're getting nowhere near the expected 300-600 t/s based on the hardware capabilities—the MoE optimization with NVFP4 marlin and cutlass kernels only managed 28 t/s, which is still far too low."

The assistant immediately acknowledges that both NVFP4 paths (marlin and cutlass) produce identical throughput, which is itself a diagnostic clue: if two fundamentally different MoE backends—one using W4A16 dequantization to BF16 tensor cores, the other using native FP4×FP4 tensor-core operations—produce the same result, then the bottleneck is upstream of the MoE computation. The MoE backend choice is irrelevant because something else is dominating the wall-clock time.

What follows is one of the most instructive passages in the entire session: a careful, step-by-step roofline analysis that exposes the true magnitude of the problem.

The Napkin Math

The assistant works through the arithmetic:

  1. Model size: 13B active parameters in FP4 (0.5 bytes per parameter) = ~6.5 GB of weights to read per decode step.
  2. TP4 sharding: Across 4 GPUs, each GPU reads ~1.6 GB of weights.
  3. Memory bandwidth: Each RTX PRO 6000 Blackwell GPU has ~1.9 TB/s of VRAM bandwidth.
  4. Theoretical minimum: 1.6 GB ÷ 1.9 TB/s ≈ 0.86 ms per GPU for weight loading.
  5. Adding KV cache and activations: Even with overhead, a single decode step should take 1–10 ms.
  6. Expected throughput at 10 ms/step: 100 steps/second × 16 tokens/step = 1,600 tok/s.
  7. Expected throughput at 1 ms/step: 16,000 tok/s. But the measured throughput is 28 tok/s. At C=16 with 256 output tokens, each step takes approximately 770 ms (16 tokens ÷ 28 tok/s × 256 tokens worth of steps... actually, let's be precise: 28 tok/s at C=16 means the system generates 28 tokens per second total. Since there are 16 concurrent requests, each step generates 16 tokens, so the step time is 16 ÷ 28 ≈ 570 ms per step. The assistant's own calculation puts it at ~770 ms per step from profile data, but either way, the gap is enormous.) The assistant states it bluntly: "That's a 200× gap from what memory bandwidth alone would allow. This isn't a memory-bound or compute-bound problem; something is catastrophically inefficient in how the decode is executing." This is the critical insight. The assistant had been operating in a "kernel optimization" frame—tuning MoE backends, adjusting CUDA graphs, tweaking NCCL settings. But the roofline analysis reveals that even the slowest possible memory-bound implementation should be orders of magnitude faster. The problem cannot be in the kernels themselves; it must be structural.

Decomposing the 770 ms

The assistant then decomposes the 770 ms per step using profile data:

"With a 13B model in FP4 across 4 GPUs, I'm touching about 1.6GB of weights per step. At 1.9TB/s bandwidth, that's roughly 0.85ms just for weight reads. The KV cache access for 16 requests with 512 context is negligible. So the memory-bound floor is around 1ms per step, which would give us 16,000 tokens/second. With realistic overhead from kernel launches and communication, I'd expect 10-30ms per step, landing us at 500-1600 tokens/second—which matches what the user is targeting."

This second calculation is even more damning. Even with generous overhead assumptions (10–30 ms per step instead of the theoretical 1 ms), the expected throughput is 500–1,600 tok/s. The measured 28 tok/s represents a 30–70× shortfall from this already-conservative estimate.

The MQA Redundancy Insight

The assistant also identifies a specific inefficiency in the attention mechanism. DeepSeek-V4-Flash uses Multi-Query Attention (MQA), where all 64 query heads share a single KV head. The assistant realizes:

"The attention kernel is doing something wasteful—with MQA where all 64 query heads share a single KV head, the kernel is re-reading the same KV data 64 times from global memory instead of sharing it. That's 13GB of redundant reads per step just for attention, which alone could account for several milliseconds."

This is a subtle but important observation. In MQA, the KV cache is shared across all query heads, but if the attention kernel is naively implemented, each query head independently loads the KV data, multiplying the memory traffic by 64. This could add ~7 ms of redundant reads (13 GB ÷ 1.9 TB/s), which is wasteful but still doesn't explain the 300 ms attention kernel time. The assistant correctly notes that "even that doesn't fully explain the 300ms bottleneck."

The Amdahl's Law Check

Perhaps the most sophisticated reasoning in the message is the assistant's application of Amdahl's Law to check its own assumptions:

"Even if I fix both the attention and MoE kernels, Amdahl's law suggests I'd only get to ~45 t/s, nowhere near 300-600 — so either the profile is missing where the wall-clock time actually goes, or the batch isn't really running at C=16."

This is a crucial sanity check. If the profile says the two dominant kernels account for 78% of the time, and even a perfect optimization of those kernels (hypothetically reducing their contribution to zero) would only yield a ~4.5× speedup (from 28 to ~45 tok/s), then the profile must be wrong or incomplete. The wall-clock time must be going somewhere that the profiler isn't capturing—perhaps gaps between kernels, scheduler overhead, or communication stalls.

This leads the assistant to the most likely explanation: the batch isn't actually running at C=16.

The Diagnostic Bash Command

The reasoning culminates in a targeted diagnostic command. The assistant stops speculating and goes directly to the source of truth: the server logs. It queries the decode batch lines during the C=16 benchmark to check the #running-req field—the actual number of requests being processed simultaneously.

The command also checks server configuration parameters like num_continuous_decode_steps, max_running_requests, disable_overlap_schedule, schedule_conservativeness, and chunked_prefill_size—all of which could affect the effective batch size.

The output is devastating:

[2026-06-17 21:09:09 TP0] Decode batch, #running-req: 8, #full token: 2560, ...
[2026-06-17 21:09:19 TP0] Decode batch, #running-req: 8, #full token: 3072, ...

#running-req: 8 — not 16. The batch size is half of what was expected.

This single line explains the entire mystery. The system is decoding batches of 8 tokens, not 16. The throughput of 28 tok/s at C=16 was actually 28 tok/s at C=8, which means the per-request throughput is even worse than it appeared. And the fact that the batch is only 8 explains why the roofline analysis showed a 100× gap: the assistant was calculating based on C=16, but the system was only delivering C=8.

But why is the batch only 8? The assistant's reasoning doesn't fully explore this in the target message itself—that exploration happens in subsequent messages. But the target message sets up the diagnostic framework that makes the discovery possible.

Assumptions and Their Consequences

This message is particularly valuable for the assumptions it exposes and corrects:

Assumption 1: The Benchmark is Running at the Requested Concurrency

The assistant had been running benchmarks with C=16 and assuming that 16 concurrent requests were being processed simultaneously. The decode logs reveal this assumption was false. The scheduler was only batching 8 requests at a time, likely due to memory constraints, scheduler conservativeness, or the MTP verifier consuming GPU memory that reduced the available batch capacity.

Assumption 2: The Profile Captures All Wall-Clock Time

The assistant's earlier profiling showed two kernels consuming 78% of GPU time, suggesting that optimizing those kernels would yield at most a ~4.5× speedup. But the Amdahl's Law check revealed that even a 4.5× speedup would only reach ~45 tok/s, not the 300–600 tok/s target. This forced the realization that the profile was missing significant wall-clock time—likely scheduler gaps, synchronization stalls, or communication overhead that the GPU profiler doesn't attribute to any kernel.

Assumption 3: Kernel Optimization is the Highest-Leverage Path

The assistant had spent days optimizing MoE backends (marlin, cutlass, flashinfer) and attention kernels. The roofline analysis shows that even perfect kernel optimization cannot close a 100× gap. The real leverage is in fixing the structural issue—the batching behavior—not in micro-optimizing individual kernels.

Assumption 4: The MoE Backend Choice Matters

Both NVFP4 backends (marlin and cutlass) produced identical throughput at C=16. This is a strong signal that the bottleneck is upstream of the MoE computation. If the MoE were the bottleneck, different backends with different computational efficiencies would produce different throughputs. Their identity suggests the MoE is waiting on something else—likely the scheduler or the attention kernel.

Input Knowledge Required

To fully understand this message, the reader needs:

  1. Roofline analysis fundamentals: The concept of memory-bound vs. compute-bound kernels, and how to calculate theoretical minimum execution time from memory bandwidth and data volume.
  2. FP4 quantization: Understanding that FP4 (4-bit floating point) stores each parameter in 0.5 bytes, so a 13B parameter model requires ~6.5 GB of storage.
  3. Tensor parallelism (TP4): The model is sharded across 4 GPUs, so each GPU holds 1/4 of the weights (~1.6 GB for the active parameters).
  4. Mixture-of-Experts (MoE): The model has 13B active parameters but many more total parameters; only the top-k experts are activated per token.
  5. Multi-Query Attention (MQA): A attention variant where multiple query heads share a single KV head, which can lead to redundant memory reads if the kernel isn't optimized.
  6. SGLang server architecture: Understanding decode batches, running requests, CUDA graphs, and scheduler parameters.
  7. Amdahl's Law: The principle that the speedup of a system is limited by the fraction of time spent in the part being improved.
  8. The history of the session: The NVFP4 checkpoint switch, the PR #25820 application, the marlin vs. cutlass comparison, and the earlier profiling results.

Output Knowledge Created

This message produces several important pieces of knowledge:

  1. The 100× gap diagnosis: A clear, quantified statement that the system is ~100× slower than roofline limits, which shifts the optimization strategy from kernel micro-tuning to structural diagnosis.
  2. The batch size discovery: The decode logs reveal #running-req: 8 instead of the expected 16, providing the first concrete evidence of a scheduling/batching problem.
  3. The MQA redundancy insight: Identification of the 64× redundant KV cache read as a potential source of inefficiency in the attention kernel.
  4. The Amdahl's Law sanity check: A demonstration of how to use Amdahl's Law to validate whether profiling data is consistent with observed wall-clock time.
  5. A diagnostic methodology: The message demonstrates a repeatable approach: (a) compute roofline bounds, (b) compare to measured performance, (c) if the gap is too large to be explained by kernel inefficiency, look for structural issues, (d) check the scheduler logs to verify the actual operating conditions.

The Thinking Process: A Window into Diagnostic Reasoning

The agent reasoning section of this message is particularly valuable because it shows the assistant's thought process in real-time, including false starts and self-corrections.

The reasoning begins with a straightforward observation: both NVFP4 paths give the same throughput, so the MoE backend isn't the lever. Then the assistant performs the roofline calculation, which produces a shocking result: the theoretical minimum step time is <1 ms, but the measured step time is ~770 ms.

At this point, the assistant cycles through possible explanations:

Significance and Implications

This message represents a turning point in the optimization campaign. Before this message, the assistant was operating in a "kernel optimization" paradigm: identify the slowest kernel, optimize it, measure the improvement. After this message, the paradigm shifts to "structural diagnosis": identify why the system isn't utilizing its hardware efficiently at the macro level.

The discovery that the batch size is only 8 (not 16) opens up an entirely new line of investigation: Why is the scheduler limiting the batch? Is it a memory constraint (the MTP verifier consuming GPU memory)? A scheduler configuration issue (conservativeness, overlap schedule)? A NUMA/cgroup limitation? Each of these questions leads to different fixes.

More broadly, this message illustrates a fundamental principle of systems optimization: always check your assumptions against reality before optimizing. The assistant spent days optimizing MoE kernels, but the real bottleneck was in the scheduler—a layer of the system that the profiling tools don't directly expose. The roofline analysis provided the necessary reality check: when theory says you should be 100× faster, you're not looking at the right problem.

The message also demonstrates the value of the user's challenge. The assistant had become comfortable with the ~28 tok/s number, accepting it as a hardware limitation. The user's insistence that the hardware should deliver 300–600 tok/s forced the re-examination that led to the real diagnosis. This is a powerful reminder that in optimization work, external benchmarks and theoretical limits are essential guardrails against premature acceptance of suboptimal performance.

Conclusion

Message 12515 is a masterclass in diagnostic reasoning under pressure. Faced with a 100× gap between roofline theory and measured performance, the assistant systematically works through the arithmetic, checks its assumptions, applies Amdahl's Law as a consistency check, and finally goes to the source of truth—the server logs—to discover that the batch size is half of what was expected.

The message is notable not just for its technical content but for its demonstration of intellectual rigor: the willingness to abandon a comfortable optimization narrative when the numbers don't add up, the discipline to compute theoretical bounds before diving into empirical tuning, and the creativity to use Amdahl's Law as a sanity check on profiling data.

For anyone working on LLM inference optimization, this message offers a template for diagnosing performance issues: start with the roofline, check your assumptions, verify the scheduler logs, and never accept a 100× gap without understanding where the time is actually going. The answer is often not in the kernels—it's in the system architecture that orchestrates them.